背景
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×2048
int 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_sample
Max error: 0.000000
まとめ
- CUDAでプログラミングをする際のテンプレートを調査、記載した
参考文献
変更履歴
- 2019/08/26: 新規作成
