1次元配列の和をGPUで計算 (CUDA vs HMPP)
CUDA
1次元配列の和を計算CPUコードをCUDA化するために、追加・修正したコードの一部を抜粋して以下に示します。
__global__ void cu_kernel(int n, float* a_gpu, float* b_gpu, float* res_gpu){
int i;
i = blockIdx.x*blockDim.x + threadIdx.x ; // GPUスレッドのインデックス
res_gpu[i] = a_gpu[i] + b_gpu[i];
}
int main(){
dim3 thread(num_th_x); dim3 block(n/num_th_x);
cudaMalloc((void**)&a_gpu,sizeof(float)*n);
cudaMemcpy(a_gpu, a, sizeof(float)*n, cudaMemcpyHostToDevice);
cu_kernel<<<block, thread>>> (n, a_gpu, b_gpu, res_gpu);
cudaMemcpy(res, res_gpu, sizeof(float)*n, cudaMemcpyDeviceToHost);
cudaFree(res_gpu);
}
HMPP
HMPPでは元のCPUコードの2箇所にディレクティブを追加するだけで自動的にループをグリッド化してGPUで計算することができます。
#pragma hmpp k1 codelet, target=CUDA &
#pragma hmpp k1 args[res].io=out
void kernel(int n, float a[n], float b[n], float res[n]){
int i;
for(i=0;i<n;i++){
res[i] = a[i] + b[i];
}
}
int main(){
・・・
// codeletの呼び出し。GPUメモリ割り当て/解放、データ転送、kernel実行が行われる
#pragma hmpp k1 callsite
kernel(n, a, b, res);
・・・
}
----------------------------------------------------------------------------------
HMPPではCPUコードの2箇所にディレクティブを追加し、他に変更を行う必要がありません。codeletで指定された関数内のループは自動的にグリッド化されます。メモリ確保やデータ転送はHMPPにより自動的に行われますが、ディレクティブによってユーザーがコントロールすることも可能です。ここでは、2箇所にディレクティブを追加するだけでCPUコードを保持してGPUで計算を行うことができており、CUDAに比べて非常に簡単です。
OpenMPの様な”#ディレクティブ”はCPUコンパイラ―では無視され、ソース保全されます。
CUDAではCPUコードにcudaMallocやcudaMemcpyなどの関数を追加してメモリ確保やデータ転送を行い、カーネルの計算では、ループをGPUのスレッドに割り当てた形に書き換えてコードを作成する必要があります。
---------------------------------------------------------------------------------- |
|
