背景
CUDAを利用して並列化を行い、処理を高速化する必要が出てきたため、CUDAで並列化処理を記述する方法について記述する
記事の目的
CUDAで並列化処理を行う際のテンプレートを作成する
CUDA
ここでは、CUDAを利用したCプログラムの記述方法について記載する。
CUDAとは
CUDAは、nvidia社が提供するGPUを利用した並列演算プログラミング基盤である利点
- ユーザーが多く、情報も入手しやすい
- ライブラリが充実している
- 導入が容易である
テンプレート
CUDAを利用して並列処理を行うテンプレートを記載する。
// cuda_sample_code.cu#include <stdio.h>// デバイス(GPU)側のスレッドの設定 //////////////////////////////////// BLOCK_SIZE_Xは、1ブロックあたりのスレッド数を表す// GPUの種類により、1ブロックあたりのスレッド数の制限が異なる// 最適なスレッド数を設定しないと、カーネル関数の処理がスキップされる// 注)上記の場合、エラーでプロセスが落ちる事はない#define BLOCK_SIZE_X 512// カーネル関数(GPUで処理する関数)vec_sumの宣言 /////////////////////// この関数がGPUで並列実行される__global__void vec_sum(float k, float *a, float *b, float *c){// iは、"for(int i=0; i < grid.x * block.x; i++)" の値を取る// ただし、並列処理のため、i++順に処理されるわけではないint i = blockIdx.x*blockDim.x + threadIdx.x;c[i] = k*a[i] + b[i];// GPU内のスレッドの同期 ////////////////////////////////////////// block_size * grid_size個の全てのスレッドが、ここで同期する__syncthreads();}// カーネル関数を呼び出す関数を宣言 /////////////////////////////////void cuda_handler(float *a, float *b, float *c){// デバイス用のポインタを宣言 ///////////////////////////////////float *d_a, *d_b, *d_c;// デバイス(GPU)側の配列のメモリを確保 //////////////////////////// デバイス側のメモリ確保処理は重いため、回数は減らした方が良いcudaMalloc(&d_a, N*sizeof(float));cudaMalloc(&d_b, N*sizeof(float));cudaMalloc(&d_c, N*sizeof(float));// ホスト側の配列内のデータをデバイス側にコピー /////////////////cudaMemcpy(d_a, a, N*sizeof(float), cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, N*sizeof(float), cudaMemcpyHostToDevice);// デバイス側の配列cの全ての要素に0をセット /////////////////////cudaMemset(d_c, 0, N*sizeof(float));// cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////// ここでは、メモリへのデータセットのエラーチェックに利用している// 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない// 注)エラーチェックをアクティブにすると、性能が極端に落ちるcheckCudaErrors(cudaGetLastError());// 並列処理するスレッド数を定義 /////////////////////////////////// 総スレッド数は block_size * grid_size = N 個である// x, y, zの3次元まで設定可能であるdim3 block_size (BLOCK_SIZE_X, 1, 1);dim3 grid_size (N / block.x, 1, 1);// カーネル関数(GPUで処理する関数)の呼び出し ////////////////////// カーネル関数内では、デバイス側のメモリ内のデータのみ操作可能vec_sum<<<grid_size, block_size>>>(2.0f, d_a, d_b, d_c);// cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////// ここでは、vec_sumのエラーチェックに利用している// 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない// 注)エラーチェックをアクティブにすると、性能が極端に落ちるcheckCudaErrors(cudaGetLastError());// この行までに実行されたカーネル関数の処理が完了するまで待機 ///// デバイスとホストの処理は非同期である// 同期処理を行うか、cudaMemcpyするまで、互いは独立して動作するcudaThreadSynchronize();// 計算結果をホスト側にコピー ///////////////////////////////////cudaMemcpy(c, d_c, N*sizeof(float), cudaMemcpyDeviceToHost);// デバイス(GPU)側の配列のメモリを開放 //////////////////////////cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);// cudaを利用した処理のうち、最後のエラーを取得し、表示 /////////// ここでは、メモリへのデータセットのエラーチェックに利用している// 注)nvccのコンパイルオプションに"-g -G"を追加しないと動作しない// 注)エラーチェックをアクティブにすると、性能が極端に落ちるcheckCudaErrors(cudaGetLastError());}int main(void){// 計算回数の設定 ///////////////////////////////////////////////// N = 512×2048int N = 1<<20;// ホスト用のポインタを宣言 /////////////////////////////////////float *a, *b, *c;// ホスト側の配列のメモリを確保 /////////////////////////////////a = (float*)malloc(N*sizeof(float));b = (float*)malloc(N*sizeof(float));c = (float*)malloc(N*sizeof(float));// a, bの配列にそれぞれ1，2を代入////////////////////////////////for (int i = 0; i < N; i++) {a[i] = 1.0f;b[i] = 2.0f;}// cudaでの処理を行う関数 ///////////////////////////////////////cuda_handler(a, b, c);// 計算結果の確認 ///////////////////////////////////////////////float maxError = 0.0f;for (int i = 0; i < N; i++){maxError = max(maxError, abs(c[i]-4.0f));}printf("Max error: %f", maxError);// ホスト側の配列のメモリを開放 /////////////////////////////////free(a);free(b);free(c);return 0;}
テンプレートのコンパイル方法
テンプレートのコンパイル方法を記載する。デバッグログを出力する場合は、下記のように記載する。
nvcc -lcuda ./cuda_sample_code.cu -o ./cuda_sample
nvcc -lcuda -g -G ./cuda_sample_code.cu -o ./cuda_sample
テンプレートの実行結果
テンプレートの実行結果を記載する。
$ ./cuda_sampleMax error: 0.000000
まとめ
- CUDAでプログラミングをする際のテンプレートを調査、記載した
参考文献
変更履歴
- 2019/08/26: 新規作成
0 件のコメント:
コメントを投稿