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ではどうなってるか未確認です。ごめんね∩(・ω・)∩