コードカキマス反撃のシゲハル

新しく学んだことをまとめます

CUDAによるGPUプログラミング

春に下書きして、投稿してませんでした。

GPUとはGraphics Processing Unit の略称で、PCやワークステーションで画像処理を担当する部品の一つです。ゲーミングPCとかだったらグラフィックボードを大体積んでますね。画面の描写の役割とかを果たしてくれます。GPUを利用することで、大量のデータを複数のプロセッサで同時かつ並列処理することができます。

今大学で勉強中のCUDAを利用した、GPUプログラミングの基本をまとめておこうかと思います。

 

基本の流れは以下です。

f:id:aspr_aspr:20170708030556p:plain

ここで、カーネルとはGPU側に担当してもらう処理関数みたいなものです。

 

具体的にどういう風に記述するのか、GPUの構造を確認しながらまとめます。

f:id:aspr_aspr:20170708031114p:plain

出典:「CUDAプログラミングの基本 パートⅠ」 by NVIDIA

GPUで関数を実行する時に特に意識するのがこの、グリッド、ブロック、スレッドです。スレッドはカーネル関数が動作させた時のプログラムの最小単位を表します。これが同時に何個も動作することで高速な処理が可能です。

ブロックは複数のスレッドの集合体で、グリッドは複数のブロックの集合体と考えれば簡単化と思います。

ブロック、スレッドは1次元や2次元などで定義することができます。

 

ここで、画像の簡単な処理を実行したいと思います。使う画像は「lenna」という画像処理分野ではよく使われる綺麗なお姉さんの画像を使って画像の二値化(白黒に)します。

ソースは以下。

 

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>


#define N 256


__global__ void shirokuro(unsigned char *original_image_gpu, unsigned char *monokuro_image_gpu) {
 int i,j;
// 以下の2行はブロック、スレッドの通し番号を表す。
 i= blockIdx.x * blockDim.x + threadIdx.x;
 j= blockIdx.y * blockDim.y + threadIdx.y;
//以下の行で、全てのスレッドで同時実行できる。
 if(original_image[i*N+j] < 128){
  monokuro_image_gpu[i*N+j] = 0
 }
 else{
  monokuro_image_gpu[i*N+J] = 255
 }
}


int main(void){
 unsigned char *original_image, *monokuro_image;
 unsigned char *original_image_gpu, *monokuro_image_gpu;
 float time_ms = 0.0f;//時間計測の変数
 cudaEvent_t start, stop;//時間計測の為のイベント
 FILE *fp;


// host側で元画像と処理後画像のメモリ領域を確保
 original_image = (unsigned char*)malloc(sizeof(unsigned char)*N*N);
 monokuro_image = (unsigned char*)malloc(sizeof(unsigned char)*N*N);


//元画像、「lennna.256」を配列に格納
 fp = fopen("lennna.256", "rb");
 fread(f, sizeof(unsigned char), N*N, fp);
 fclose(fp);


//cudaにおける時間計測イベント作成
 cudaEventCreate(&start);
 cudaEventCreate(&stop);


// cuda側(GPU)でメモリ確保
 cudaMalloc((void**)&original_image_gpu, sizeof(unsigned char)*N*N);
 cudaMalloc((void**)&monokuro_image_gpu, sizeof(unsigned char)*N*N);


// cuda側に元画像をコピー
 cudaMemcpy(original_image_gpu, original_image, sizeof(unsigned char)*N*N, cudaMemcpyHostToDevice);


// block, threads の確保
 dim3 blocks(16,16,1);
 dim3 threads(16,16,1);


// cuda時間計測イベントで計測開始
 cudaEventRecord(start, 0);


//カーネル関数の実行
 shirokuro<<< blocks, threads >>> (original_image_gpu, monokuro_image_gpu);


//実行が終わり、monokuro_image_gpu に処理後の画像が入っているので、それをmonokuro_imageに渡す
 cudaMemcpy(monokuro_image, monokuro_image_gpu, sizeof(unsigned char)*N*N, cudaMemcpyDeviceToHost);


//GPUメモリ解放
 cudaFree(monokuro_image_gpu);
 cudaFree(original_image_gpu);


// cpu側でメモリ解放
 free(original_image);
 free(monokuro_image);


// ホスト側の配列から画像生成
 fp = fopen("lenna_new.256", "wb");
 fwrite(monokuro_image, sizeof(unsigned char), N*N, fp);
 fclose(fp);


//cuda時間計測終了
 cudaEventRecord(stop, 0);
 cudaEventSynchronize(stop);


//かかった時間を出力
 cudaEventElapsedTime(&time_ms, start, stop);
 printf("time: %f ms\n",time_ms);


//不要になったcudaの時間計測イベントを消去
 cudaEventDestroy(start);
 cudaEventDestroy(stop);
}