【CUDA】グローバルメモリを2次元で使う

CUDAのスレッドは２次元で持つことができるので、その使い方の解説をします。

以下に紹介する処理で、スレッドを２次元で利用するコーディングの例を示します。

処理内容
コーディング例

処理内容

別の記事でも紹介しましたが、以下のような処理を作成します。

・ 8000×8000のint型二次元配列で、ある要素を中心として、その周辺３ｘ３の平均値を求める。
・結果を別の二次元配列に保存する。
・この処理を全要素に行う。

コーディング例

カーネルを呼び出す側の関数になります。


//処理関数
void Process_MeanGPU()
{
	int sizex = 8192;
	int sizey = 8192;
	int length = sizex * sizey;
	size_t size = sizeof(int) * length;


	//ホストメモリのポインタ
	int* pHostSrc;		//数値
	int* pHostResult;	//結果
	//デバイスメモリのポインタ
	int* pDevSrc;		//数値
	int* pDevResult;	//結果


	//ホストメモリの確保
	cudaMallocHost(&pHostSrc, size);
	cudaMallocHost(&pHostResult, size);
	//デバイスメモリの確保
	cudaMalloc(&pDevSrc, size);
	cudaMalloc(&pDevResult, size);

	//ホストメモリに適当な値を設定
        unsigned int value = 0 ;
	for (int n = 0; n < length; n++) {
		pHostSrc[n] = value++ % 1000;
	}

	//ホスト->デバイスへ数値を転送
	cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

	//カーネル
	dim3 block(128, 1, 1);
	dim3 grid((sizex + 128 - 1) / 128, sizey, 1);
	kernel_mean <<<grid, block >>> (pDevSrc, pDevResult, sizex, sizey); 

	//デバイス->ホストへ結果を転送
	cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

	//デバイスメモリの開放
	cudaFree(pDevSrc);
	cudaFree(pDevResult);
	//ホストメモリの開放
	cudaFreeHost(pHostSrc);
	cudaFreeHost(pHostResult);
}

//処理関数

void Process_MeanGPU()

{

int sizex = 8192;

int sizey = 8192;

int length = sizex * sizey;

size_t size = sizeof(int) * length;

//ホストメモリのポインタ

int* pHostSrc; //数値

int* pHostResult; //結果

//デバイスメモリのポインタ

int* pDevSrc; //数値

int* pDevResult; //結果

//ホストメモリの確保

cudaMallocHost(&pHostSrc, size);

cudaMallocHost(&pHostResult, size);

//デバイスメモリの確保

cudaMalloc(&pDevSrc, size);

cudaMalloc(&pDevResult, size);

//ホストメモリに適当な値を設定

unsigned int value = 0 ;

for (int n = 0; n < length; n++) {

pHostSrc[n] = value++ % 1000;

}

//ホスト->デバイスへ数値を転送

cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

//カーネル

dim3 block(128, 1, 1);

dim3 grid((sizex + 128 - 1) / 128, sizey, 1);

kernel_mean <<<grid, block >>> (pDevSrc, pDevResult, sizex, sizey);

//デバイス->ホストへ結果を転送

cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

//デバイスメモリの開放

cudaFree(pDevSrc);

cudaFree(pDevResult);

//ホストメモリの開放

cudaFreeHost(pHostSrc);

cudaFreeHost(pHostResult);

}

こちらはカーネルです。

//平均を取得するカーネル
__global__ void kernel_mean(int* pSrc, int *pResult, int sizex, int sizey)
{
	int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex
	int idy = blockDim.y * blockIdx.y + threadIdx.y; //自分のスレッドyのindex

	if (idx < 1 || idy < 1 || idx >= sizex - 1 || idy >= sizey - 1) {
		return; //配列の左右上下端は無視
	}

	unsigned int sum = 0;

	for (int y = 0; y < 3; y++)
	{
		for (int x = 0; x < 3; x++) {
			sum += pSrc[(idy + y - 1) * sizex + (idx + x - 1)];
		}
	}

	pResult[idy * sizex + idx] =  sum / 9;

	return;
}

//平均を取得するカーネル

__global__ void kernel_mean(int* pSrc, int *pResult, int sizex, int sizey)

{

int idx = blockDim.x * blockIdx.x + threadIdx.x; //自分のスレッドxのindex

int idy = blockDim.y * blockIdx.y + threadIdx.y; //自分のスレッドyのindex

if (idx < 1 || idy < 1 || idx >= sizex - 1 || idy >= sizey - 1) {

return; //配列の左右上下端は無視

}

unsigned int sum = 0;

for (int y = 0; y < 3; y++)

{

for (int x = 0; x < 3; x++) {

sum += pSrc[(idy + y - 1) * sizex + (idx + x - 1)];

}

pResult[idy * sizex + idx] = sum / 9;

return;

}

int idx = blockDim.x * blockIdx.x + threadIdx.x;
int idy = blockDim.y * blockIdx.y + threadIdx.y;

の記述で、自分のスレッドのx,yインデックスを計算しています。
おまじないのようなものなので覚えておきましょう。