__device__なclassを作るとどうなるかというお話。
普通に使えたらかなり便利じゃね？
とりあえず、
- __global__ void
- test(int *value)
- {
- value[0] = hoge.func() - hoge.func2();
- }
なんてのを用意しておいてみる。
- class CHoge
- {
- public:
- int value;
- CHoge(){
- value = 99;
- }
- ~CHoge(){
- }
- int func()
- {
- return 98;
- }
- int func2()
- {
- return value;
- }
- }hoge;
でどうか。
結果は、hoge.func()とhoge.func2()の両方で
"cuda.cu", line 78: error: calling a host function from a
__device__/__global__ function is only allowed in device emulation
mode
value[0] = hoge.func() - hoge.func2();//value;
^
"cuda.cu", line 78: error: calling a host function from a
__device__/__global__ function is only allowed in device emulation
mode
value[0] = hoge.func() - hoge.func2();//value;
^
2 errors detected in the compilation of "/tmp/tmpxft_00007432_00000000-3.ii".
となり、関数が呼べない。まぁそりゃそうだろう。
ちょっと問題を縮小し、最低限の記述に落としてみる。
- class CHoge
- {
- public:
- int func()
- {
- return 98;
- }
- int func2()
- {
- return 99;
- }
- }hoge;
これでも同じエラー。
- class CHoge
- {
- public:
- __device__ int func()
- {
- return 98;
- }
- __device__ int func2()
- {
- return 99;
- }
- }hoge;
これだとうまく動いて-1が得られる。なるほど。
ちょっと趣向を変えて。
- class CHoge
- {
- public:
- __device__ int func()
- {
- return 98;
- }
- __device__ int func2()
- {
- return 99;
- }
- };
- class CHoge hoge1;
- class CHoge hoge2;
- __global__ void
- test(int *value)
- {
- value[0] = hoge1.func() - hoge2.func2();
- }
- 1。ふむ。
……あれ？classの実体宣言には__device__つけていないんだけど、いいのかなぁ？
classに変数を持たせてみる。
- class CHoge
- {
- public:
- int value;
- __device__ int set(int n)
- {
- value = n;
- return value;
- }
- __device__ int func()
- {
- return 1;
- }
- __device__ int func2()
- {
- return value;
- }
- };
- class CHoge hoge1;
- class CHoge hoge2;
- __global__ void
- test(int *value)
- {
- hoge1.set(10);
- hoge2.set(20);
- value[0] = hoge1.func2() + hoge2.func2();
- }
すると、
"cuda.cu", line 62: error: identifier "hoge1" is undefined
(hoge1.value) = 10;
^
"cuda.cu", line 63: error: identifier "hoge2" is undefined
(hoge2.value) = 20;
^
"cuda.cu", line 64: error: identifier "hoge1" is undefined
(value[0]) = (((hoge1.value)) + ((hoge2.value)));
^
"cuda.cu", line 64: error: identifier "hoge2" is undefined
(value[0]) = (((hoge1.value)) + ((hoge2.value)));
^
4 errors detected in the compilation of "/tmp/tmpxft_000074c1_00000000-4.i".
……あれ？
- __device__ class CHoge hoge1;
- __device__ class CHoge hoge2;
に変更したら、今度は
"cuda.cu", line 62: warning: expression has no effect
(hoge1.value; }
^
"cuda.cu", line 63: warning: expression has no effect
(hoge2.value; }
^
というwarnigのみになり、実行結果も30と正しそうな感じに。
まぁなんだ、要するにインライン展開されまくりなんだろうな、きっと。
とりあえず関数も持てる構造体、程度には使えると。
（ええと、C++の場合は構造体に関数持たせられるんだっけ？そういえば。）
次。
class CHoge
にコンストラクタやらデストラクタやらを追加してみる。
- CHoge()
- {
- value = 1;
- }
これで
"/usr/lib/gcc/i386-redhat-linux/4.1.1/../../../../include/c++/4.1.1/cstdlib", line 178: internal error:
can't generate code for non empty constructors or destructors on
device
div(long long __n, long long __d)
^
1 catastrophic error detected in the compilation of "/tmp/tmpxft_00007555_00000000-3.ii".
Compilation aborted.
意味がわからんエラーが出る。まぁコンストラクタやデストラクタはダメだよって言われてるわけだな。
残念ながらコンストラクタを__device__にしても状況は変わらないし、int value;を__device__にしても
"cuda.cu", line 45: error: memory qualifier on data member is not allowed
__attribute__((__device__)) int value;
^
とかエラーが変わるのみ。
うーん。
流石にまともにclassらしくclassを活用することはできなさそうだ。ちょっと残念。
それにコンストラクタとかが使えないのは微妙に不便だ。
コンストラクタっぽいことをする妥当なアイディアとしては、classを__device__に持っていることが前提だけど、初期化用の__device__関数を利用するのがいいのかな。
- class CHoge
- {
- public:
- int value;
- __device__ int set(int n)
- {
- value = n;
- return value;
- }
- __device__ int func()
- {
- return 1;
- }
- __device__ int func2()
- {
- return value;
- }
- };
- __device__ class CHoge hoge1;
- __device__ class CHoge hoge2;
- __global__ void init()
- {
- hoge1.set(10);
- hoge2.set(20);
- }
- __global__ void test(int *value)
- {
- value[0] = hoge1.func2() + hoge2.func2();
- }
- 以下、ホスト側
- // 変数の初期化のために一度GPUを動かしてみる
- dim3 threads0(1, 1, 1);
- dim3 grid0(1, 1, 1);
- init<<<grid0, threads0>>>();
- // 本番？の実行
- dim3 threads(THREAD_X, 1, 1);
- dim3 grid(BLOCK_X, 1, 1);
- test<<<grid, threads>>>(d_sd1);
それにしてもなんだこのグダグダのエントリ。まとまりわりい。