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