クラス内カーネル

CUDA 4.0の話です。古くてごめんなさい。

nvccは限定的ながらC++に対応してるということで、例えばクラスメンバがvirtualだったりするとダメとかみたいな制限はあるんですが、カーネル自体がクラスメンバになれないとは。

class ACUDAClass
{
	__global__ void kernel();
public:
	void func()
	{
		kernel<<<1, 1>>>();
	}
};

void ACUDAClass::kernel()
{
}

int main()
{
	ACUDAClass ob;
	ob.func();
}

こんなエラーが出ます。

% nvcc prog.cu
prog.cu(3): error: illegal combination of memory qualifiers

1 error detected in the compilation of "/tmp/tmpxft_00003d82_00000000-4_prog.cpp1.ii".

もちろん、こういう風にクラスメンバからカーネルを呼ぶのは何も問題ありません。

__global__ void kernel()
{
}

class ACUDAClass
{
public:
	void func()
	{
		kernel<<<1, 1>>>();
	}
};

int main()
{
	ACUDAClass ob;
	ob.func();
}

理由はどうも単純なもので、thisをカーネルに渡せないからクラスメンバとしては置けないよ!!というアレみたい。

ならばと思ってカーネルをstaticにしても同じエラーがでます。はて。
例えばグローバル変数(ホスト側にある)を参照するとundefined symbolってちゃんとエラー吐けてるから同じ程度の問題のような気もする。

ところが、__device__だとおっけーなんですね。
ふむ。わからん。

誰かホントのところ教えてくだしあ。
あと最初に言ったようにこれCUDA 4.0での話なので4.1ではどうなってるか未確認です。ごめんね∩(・ω・)∩

コメントを残す

メールアドレスが公開されることはありません。