最終更新: mikk_ni3_92 2010年03月03日(水) 15:51:46履歴
現在地 >> メニュー >> CUDA >> CUDA編02
INDEX: CUDA編01 << CUDA編02 >> CUDA編03
「ローカルメモリ」は「Thread」が使うメモリ。
「共有メモリ」は「Block」内で使うメモリ。つまりそのブロック内の「Thread」で共有される。
「グローバルメモリ」はGPU全体で使用できるメモリ。
GPUへのデータ転送はcudaMemcpy関数を使う
【例】:GPUへのデータ転送
cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
「count」バイト分だけ「src」の内容を「dst」へコピー
「kind」には
CPUでGPUが実行する関数をコールすることでGPUでの処理がはじまる。
※GPUでの関数の事をカーネルともよぶそうだ
【例】:カーネルの呼び出し
詳細については
ここまでをまとめると次のようになる
GPUへ渡した配列データを処理するにはなんとかして、アクセスするための添え字を取得する必要がある。
CUDAでは組み込み変数として、threadIdxやblockDimなどがある。
これらを使って、配列データにアクセスする。
詳細は
【例】
まず、現在のスレッド番号、ブロック番号から配列データのどこに対応するかを計算。
処理すべき配列番号を取得したらそこにアクセスして、何らかの処理をおこなう。
void __syncthreads()
ブロック内のスレッドがすべてその関数がかかれた所にくるまで待つ。
ブロック内で同期処理をする時に使う。
INDEX: CUDA編01 << CUDA編02 >> CUDA編03
「ローカルメモリ」は「Thread」が使うメモリ。
「共有メモリ」は「Block」内で使うメモリ。つまりそのブロック内の「Thread」で共有される。
「グローバルメモリ」はGPU全体で使用できるメモリ。
GPUにデータを転送するためには、まずGPU上にメモリを確保する必要がある。
【例】:GPU上にメモリ確保
cudaError_t cudaMalloc (void **devPtr, size_t size)
GPU上にメモリを確保する。
※CUDA_SAFE_CALL()でエラーを拾う。(自分でチェックしてもよい)
不要になったら
cudaError_t cudaFree (void *devPtr)
でメモリを破棄する。
【例】:GPU上にメモリ確保
const int AryNum = 20000; … … //GPUのメモリ確保 int DataSize = sizeof(float)*AryNum; float *gpuAry; CUDA_SAFE_CALL(cudaMalloc( reinterpret_cast<void**>( &gpuAry), DataSize) ); … … cudaFree(gpuAry);//メモリ破棄【解説】
cudaError_t cudaMalloc (void **devPtr, size_t size)
GPU上にメモリを確保する。
※CUDA_SAFE_CALL()でエラーを拾う。(自分でチェックしてもよい)
不要になったら
cudaError_t cudaFree (void *devPtr)
でメモリを破棄する。
GPUへのデータ転送はcudaMemcpy関数を使う
【例】:GPUへのデータ転送
//CPUでメモリ確保 float *ary = new float[AryNum]; … 値の設定 … //GPUのメモリ確保 int DataSize = sizeof(float)*AryNum; float *gpuAry; … GPUのメモリ確保 … CUDA_SAFE_CALL( cudaMemcpy( gpuAry, ary, DataSize , cudaMemcpyHostToDevice));【解説】
cudaError_t cudaMemcpy (void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)
「count」バイト分だけ「src」の内容を「dst」へコピー
「kind」には
- cudaMemcpyHostToHost
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
CPUでGPUが実行する関数をコールすることでGPUでの処理がはじまる。
※GPUでの関数の事をカーネルともよぶそうだ
【例】:カーネルの呼び出し
//------------- CUDAでの処理 ---------------// __global__ void cuAddValue(float *gpuAryData) { const int threadId = blockIdx.x * blockDim.x + threadIdx.x; //スレッドIdを取得 gpuAryData[threadId] += 20000;//すべての要素を+20000する __syncthreads(); //同期をとる } … … //GPU用にブロックとスレッドを用意 dim3 grid(100,1,1);//100×1サイズのグリッド(1グリッドあたり100ブロック) dim3 block(200,1,1);//200×1のブロックサイズ(1ブロックあたり200スレッド) //グリッド100個、ブロック200 cuAddValue<<<grid,block>>>(gpuAry);
詳細については
「cudaMemcpy関数」の第4引数「にcudaMemcpyDeviceToHost」を指定すればよい。
【例】:GPUからデータを取り出す
【例】:GPUからデータを取り出す
//CPUでメモリ確保 float *ary = new float[AryNum]; … 値の設定 … //GPUのメモリ確保 int DataSize = sizeof(float)*AryNum; float *gpuAry; … GPUのメモリ確保 … … … //GPUから戻す(GPU -> Host) CUDA_SAFE_CALL( cudaMemcpy( ary,gpuAry ,DataSize, cudaMemcpyDeviceToHost) );
ここまでをまとめると次のようになる
const int AryNum = 20000; //------------- CUDAでの処理 ---------------// __global__ void cuAddValue(float *gpuAryData) { … … } //------------- メイン関数 --------------------// int main(int argc, char **argv) { CUT_DEVICE_INIT(argc,argv); //CPUでメモリ確保+値設定 float *ary = new float[AryNum]; for(int loop = 0;loop < AryNum; ++loop) { ary[loop] = loop; } //GPUのメモリ確保 int DataSize = sizeof(float)*AryNum; float *gpuAry; CUDA_SAFE_CALL(cudaMalloc( reinterpret_cast<void**>( &gpuAry), DataSize) ); //GPUへデータをコピー(Host -> GPU) CUDA_SAFE_CALL( cudaMemcpy( gpuAry, ary, DataSize , cudaMemcpyHostToDevice)); //GPU用にブロックとスレッドを用意 dim3 grid(100,1,1);//100×1サイズのグリッド(1グリッドあたり100ブロック) dim3 block(200,1,1);//200×1のブロックサイズ(1ブロックあたり200スレッド) cuAddValue<<<grid,block>>>(gpuAry); //GPUから戻す(GPU -> Host) CUDA_SAFE_CALL( cudaMemcpy( ary,gpuAry ,DataSize, cudaMemcpyDeviceToHost) ); cudaFree(gpuAry);//メモリ破棄 //出力テスト for(int loop = 0;loop < AryNum; ++loop) { std::cout << ary[loop] << "\n"; } delete [] ary;//メモリクリア CUT_EXIT(argc, argv);//終了 return 0; }
GPUへ渡した配列データを処理するにはなんとかして、アクセスするための添え字を取得する必要がある。
CUDAでは組み込み変数として、threadIdxやblockDimなどがある。
これらを使って、配列データにアクセスする。
詳細は
【例】
//------------- CUDAでの処理 ---------------// __global__ void cuAddValue(float *gpuAryData) { const int threadId = blockIdx.x * blockDim.x + threadIdx.x; //スレッドIdを取得 gpuAryData[threadId] += 20000;//すべての要素を+20000する __syncthreads(); //同期をとる } … … cuAddValue<<<grid,block>>>(gpuAry);【解説】
まず、現在のスレッド番号、ブロック番号から配列データのどこに対応するかを計算。
処理すべき配列番号を取得したらそこにアクセスして、何らかの処理をおこなう。
void __syncthreads()
ブロック内のスレッドがすべてその関数がかかれた所にくるまで待つ。
ブロック内で同期処理をする時に使う。