ひびろぐ ver.2

http://developer.nvidia.com/object/cuda.html]から開発に必要なものを入手する。具体的にはCUDA ToolkitとCUDA SDK。32bitとか書いてあるけど気にしない。
ToolkitとSDKを導入。特に文句は言われない。
- Toolkitは C:\CUDA に導入されたのを確認。
- SDKは C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK に導入されたのを確認。
Tookit導入時にご丁寧に　C:\CUDA\bin　にパスが通されていることに気がつく。いきなりnvcc.exeを実行可能。
適当なサンプルCUDAプログラムを準備。今回は自前の超単純なプログラムを利用することにする。
nvccでコンパイルしてみるが、エラー。ヘッダファイルが見つからない旨を伝えられる。
- nvcc cuda.cu
エラーメッセージを参考に、Program Files以下のMicrosoft Visual Studio 8\VC\includeとNVIDIA Corporation\NVIDIA CUDA SDK\common\incをインクルードパスとして指定することにする。
- nvcc cuda.cu -I”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\include” -I”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\inc”
ヘッダファイルが見つからない旨は消えたが、clが無いといわれる。
clと言ったらVCのアレなので、C:\Program Files (x86)\Microsoft Visual Studio 8\VC\binにパスを通してみる。
今度はmspdb80.dllが見つからないと怒られる。適当にファイル検索し、C:\Program Files (x86)\Microsoft Visual Studio 8\Common7\IDEに見つけたのでパスを通す。
LINKエラー、libcpmt.libが開けないらしい。nvcc実行時引数に-Lオプションか。
C:\Program Files (x86)\Microsoft Visual Studio 8\VC\libにターゲットを捕捉。-Lオプションを追加する。
- nvcc cuda.cu -I”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\include” -I”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\inc” -L”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\lib”
__imp__cutCheckCmdLineFlag@12の未解決エラー。ライブラリのリンク不足なのは自明。名前からして恐らくcutilライブラリ。
C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\libにcutil32.libを発見。-Lと-lのオプションを追加する。
- nvcc cuda.cu -I”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\include” -I”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\inc” -L”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\lib” -L”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\lib” -lcutil32
nvcc成功。しかし実行すると今度はcutil32.dllが無いと怒られる。ああそうか、dll使ってるのか。パス通ってないんだな。
C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\libにdllを捕捉、パスを通す。
実行。There is no device supporting CUDA.で止まる。そりゃそうだ、Radeonだし。
デバイスエミュレーションに望みを託す。nvccに-deviceemuオプションをつけてやればOK。
- nvcc cuda.cu -I”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\include” -I”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\inc” -L”C:\Program Files (x86)\Microsoft Visual Studio 8\VC\lib” -L”C:\Program Files (x86)\NVIDIA Corporation\NVIDIA CUDA SDK\common\lib” -lcutil32 -deviceemu
実行。ちゃんと動く。
- ちなみに-deviceemuの有無でバイナリサイズは変わらず。動的リンクだからでいいのかな。

というわけで無事実行に成功。

DirectXのREFみたいにnativeとemulationでできることできないことに差が生じる可能性は否定しないけど、その辺は追々ということで。

CUDAでのdeviceなclassについて

Published 10月 11th, 2007 in CUDA | コメント（0） »

__device__なclassを作るとどうなるかというお話。

普通に使えたらかなり便利じゃね？

とりあえず、

PLAIN TEXT

CUDA:

__global__ void

test(int *value)

{

value[0] = hoge.func() - hoge.func2();

}

なんてのを用意しておいてみる。

PLAIN TEXT

CUDA:

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".

となり、関数が呼べない。まぁそりゃそうだろう。

ちょっと問題を縮小し、最低限の記述に落としてみる。

PLAIN TEXT

CUDA:

class CHoge

{

public:

int func()

{

return 98;

}

int func2()

{

return 99;

}

}hoge;

これでも同じエラー。

PLAIN TEXT

CUDA:

class CHoge

{

public:

__device__ int func()

{

return 98;

}

__device__ int func2()

{

return 99;

}

}hoge;

これだとうまく動いて-1が得られる。なるほど。

ちょっと趣向を変えて。

PLAIN TEXT

CUDA:

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に変数を持たせてみる。

PLAIN TEXT

CUDA:

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".

……あれ？

PLAIN TEXT

CUDA:

__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

にコンストラクタやらデストラクタやらを追加してみる。

PLAIN TEXT

CUDA:

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__関数を利用するのがいいのかな。

PLAIN TEXT

CUDA:

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);

それにしてもなんだこのグダグダのエントリ。まとまりわりい。

CUDAコンパイラ（ｒｙ　補足

Published 10月 10th, 2007 in CUDA | コメント（0） »

変数idを__shared__にすると思い通りに動くよ！

　い　み　が　。

CUDAコンパイラの最適化が腐っている　のか？

Published 10月 10th, 2007 in CUDA | コメント（0） »

アレな状況に陥ったので報告。本家forumで議論してきたほうがいいかもシレン。

PLAIN TEXT

CUDA:

__global__ void

test(int *value)

{

int id;

id = blockIdx.x;

if(id==0){

value[0] = 1;

}else if(id==1){

value[1] = 2;

}

}

grid(2,1,1)な実行後のvalueは何になるか。

答えは

value[0] = 0

value[1] = 0

……いや、それはおかしいだろう。

実は、例えば以下のように書くと想定どおりの答えが得られる。

PLAIN TEXT

CUDA:

__global__ void

test(int *value)

{

int id;

id = blockIdx.x;

if(id==0){

value[0] = 1;

}else if(id==1){

value[id] = 2;

}

}

これで答えは

value[0] = 1

value[1] = 2

となる。

PLAIN TEXT

CUDA:

__global__ void

test(int *value)

{

int id;

id = blockIdx.x;

if(id==0){

value[id] = 1;

}else if(id==1){

value[1] = 2;

}

}

でもいいや。

PLAIN TEXT

CUDA:

__global__ void

test(int *value)

{

int id;

id = blockIdx.x;

if(id==0){

value[0] = id;

}else if(id==1){

value[1] = 2;

}

}

なんてのもありらしい。

あくまで予想なんだけど、ifブロックの中で変数idが使われていないという理由で、コンパイラがid変数を最適化の課程で殺しているんじゃないだろうか。だとしたらコンパイラのバグだよなぁ……。

やっぱ本家にゴルァかな？勘違いかもしれないので一日熟成させよう。

block間での同期を取る3

Published 10月 10th, 2007 in CUDA | コメント（0） »

このへんの続き。

前回は何故か__syncthreads()を使ったりと妙にキモイことをやっていたので、追実験。

atomic関数の存在も念頭に置いてみる。

まぁこんなケースを考えることにする。

PLAIN TEXT

CUDA:

#define BLOCK_X 8 // grid(8,1,1) を意味する

__device__ int waitflag[8] = {0,0,0,0,0,0,0,0};

__global__ void gpu

(int *mem0, int *mem1, unsigned int *_buf)

{

int id = blockIdx.x;

// 実行時間をずらすための処理

if(id==3){

mem1[0]=0;

mem1[1]=0;

mem1[2]=0;

}

// 各blockで_buf[0]を叩く

atomicAdd(&_buf[0], id);

// このへんで同期を取れないと、計算結果がblock毎に狂う可能性がある。

// 具体的には、id3以外のblockでは

// id3のatomicAddが終わる前に_buf[0]を読んでしまう可能性が高い。

// 恐らく、mem0は 26 26 26 29 26 26 26 26 になる（各要素ともデフォルトが1なので）。

// 各blockからみえる計算結果を、

// global変数に書き戻して確認

mem0[id] = _buf[0];

return;

}

んで同期をどうやるか。

とりあえず考えたのが以下の実装。

PLAIN TEXT

CUDA:

atomicAdd(&waitflag[0], 1);

while(waitflag[0]<BLOCK_X);

atomicAddで加算の保証が出来るので、あとは全blockが加算するのを待てばいいんじゃないかと。

うん。失敗した。無限ループっぽい。

直感的にはこれで行けそうなんだけど、何かひっかかるらしい。

PLAIN TEXT

CUDA:

atomicAdd(&waitflag[0], 1);

while(waitflag[0]<BLOCK_X)__syncthreads();

という感じに__syncthreads();を入れると正しく動作する。

__syncthread()にそれらしい仕様はない気がするんだが……？

ところで、atomicAddは確実に読み出して加算ができる、ってのはいいんだけど、その処理の途中で別のblock/threadが値を読んでしまうことってあるんだろうか。いや、別に「0が読めるか1が読めるかわからないよ」ってのは構わないというか当然なんだけど、何故か全然違う値が読めたりすると危険だよなぁ……。

さて、胡散臭い__syncthreads()を使わずにちゃんとやるにはどうすればいいかなんだが、作業用の変数を追加して

PLAIN TEXT

CODE:

__device__ int waitflag[8] = {0,0,0,0,0,0,0,0}; // 関数外宣言

int wait = 0; // 関数内宣言

// 計算終了通知

waitflag[id] = 1;

// 全blockが計算終了に到達していることを確認する

while(wait <BLOCK_X){

wait = 0;

for(int i=0; i<BLOCK_X; i++){

wait += waitflag[i];

}

}

とやったらうまくいった。

安全を確保するならwaitflagの変更はatomicの方がいいのかもしれないけど、まぁ大丈夫だろう。

さて。以上の振る舞いから、今回のglobal memory spaceの変数の参照はコンパイラによって最適化されてしまっているのではないかという仮説が立てられる。

nvccの-ptxや-keepで中間コードが覗けるんだけど、これ見て判断するのはちょっとツライんだよなぁ。

最後。

仮設を踏まえて適当にうまく行く記述を考えてみた。

PLAIN TEXT

CUDA:

atomicAdd(&waitflag[0], 1);

while(wait<BLOCK_X){wait=waitflag[0];}

明示的にglobal memory spaceからの再読み込みを行うような感じ。

実際、これで思い通りの動作は出来ている。

ちなみに、どの実装にしても複数回の実行にはちょっとだけ工夫が必要かもしれない。

作業用変数のクリアが必要だから。

とはいえ、作業用変数を多重化して、同期が終了するときに今使ったものとは別の作業用変数を初期化してやれば動きそうだから多分大丈夫だろう。たぶん。

以下、#ifで切ったプログラム全体。長いけど参考にベタッと貼っておく。

PLAIN TEXT

CUDA:

// -*- C++ -*-

/*

block間の同期を行ってみるテスト

*/

#include <stdlib.h>

#include <stdio.h>

#include <string.h>

#include <math.h>

#include <unistd.h>

#include <cutil.h>

#define N 8

#define BLOCK_X N

#define THREAD_X 1

//extern __device__ int* buf;

__device__ int waitflag[8] = {0,0,0,0,0,0,0,0};

__global__ void gpu

(int *mem0, int *mem1, unsigned int *_buf)

{

// -------- -------- -------- --------

int wait = 0;

int id = blockIdx.x;

// 実行時間をずらすための処理

if(id==3){

mem1[0]=0;

mem1[1]=0;

mem1[2]=0;

}

// 各blockで_buf[0]を叩く

atomicAdd(&_buf[0], id);

#if SYNC==3

// 計算終了通知

waitflag[id] = 1;

// 全blockが計算終了に到達していることを確認する

while(wait <BLOCK_X){

wait = 0;

for(int i=0; i<BLOCK_X; i++){

wait += waitflag[i];

}

}

#elif SYNC==4

atomicAdd(&waitflag[0], 1);

while(wait<BLOCK_X){wait=waitflag[0];}

#elif SYNC==2

atomicAdd(&waitflag[0], 1);

while(waitflag[0]<BLOCK_X)__syncthreads();

#else // NG

atomicAdd(&waitflag[0], 1);

while(waitflag[0]<BLOCK_X);//__syncthreads();

#endif

// ここから、各blockが計算結果を確実に取得できる

// global変数に書き戻して確認

mem0[id] = _buf[0];

// 待ち合わせデータのチェックのため、

// global変数に書き戻して確認

mem1[id] = wait;

return;

// -------- -------- -------- --------

}

// -------- -------- -------- -------- -------- -------- -------- --------

void

runTest(int argc, char** argv)

{

CUT_DEVICE_INIT();

int* h_mem0;

int* h_mem1;

int* d_mem0;

int* d_mem1;

h_mem0 = (int*)malloc(sizeof(int)*N);

h_mem1 = (int*)malloc(sizeof(int)*N);

CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem0, sizeof(int)*N));

CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem1, sizeof(int)*N));

//CUDA_SAFE_CALL(cudaMemcpy(d_mem0, h_mem0, sizeof(int)*N, cudaMemcpyHostToDevice) );

//CUDA_SAFE_CALL(cudaMemcpy(d_mem1, h_mem1, sizeof(int)*N, cudaMemcpyHostToDevice) );

unsigned int* h_buf;

unsigned int* d_buf;

CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*N));

int n;

Archive for the ‘CUDA’