【CUDA】シェアードメモリの使用方法

GPUのメモリの種類で、シェアードメモリ(SharedMemory/共有メモリ)というものがあります。

シェアードメモリは、グローバルメモリと比べて非常に高速なので、繰り返しグローバルメモリにアクセスする必要があるときは、シェアードメモリをいったん介したほうが高速化できる可能性があります。

最適化をするなら必須の要素なので、ぜひ使い方を覚えましょう。

シェアードメモリの特徴
テスト処理内容
コード例
1. 3×3 平均　シェアード無し
2. 3×3 平均　シェアード有り
速度比較
補足
1. 占有率
2. シェアードメモリの有効性

シェアードメモリの特徴

・同じブロック内のスレッドどうしなら値を共有できる。
・グローバルメモリにアクセスするより数十倍高速。
・１ブロックごとに48～64KBまで確保できる。

テスト処理内容

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

コード例

3×3 平均　シェアード無し

カーネル呼び出し部


void Process_MeanNoSharedGPU_3x3(unsigned char* pHostSrc, 
unsigned char* pHostResult, int sizex, int sizey, size_t size)
{

	//デバイスメモリのポインタ
	unsigned char* pDevSrc;		//処理前データ
	unsigned char* pDevResult;	//結果

	//ホストメモリをピンロック
	cudaHostRegister(pHostSrc, size, cudaHostRegisterDefault);
	cudaHostRegister(pHostResult, size, cudaHostRegisterDefault);
	//デバイスメモリの確保
	cudaMalloc(&pDevSrc, size);
	cudaMalloc(&pDevResult, size);
	
	//ホスト->デバイスへ数値を転送
	cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

	//カーネル
	dim3 block(128, 1);
	dim3 grid((sizex + 128 - 1) / 128, sizey);

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

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

	//デバイスメモリの開放
	cudaFree(pDevSrc);
	cudaFree(pDevResult);
	//ホストメモリのロック開放
	cudaHostUnregister(pHostSrc);
	cudaHostUnregister(pHostResult);

}

void Process_MeanNoSharedGPU_3x3(unsigned char* pHostSrc,

unsigned char* pHostResult, int sizex, int sizey, size_t size)

{

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

unsigned char* pDevSrc; //処理前データ

unsigned char* pDevResult; //結果

//ホストメモリをピンロック

cudaHostRegister(pHostSrc, size, cudaHostRegisterDefault);

cudaHostRegister(pHostResult, size, cudaHostRegisterDefault);

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

cudaMalloc(&pDevSrc, size);

cudaMalloc(&pDevResult, size);

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

cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

//カーネル

dim3 block(128, 1);

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

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

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

cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

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

cudaFree(pDevSrc);

cudaFree(pDevResult);

//ホストメモリのロック開放

cudaHostUnregister(pHostSrc);

cudaHostUnregister(pHostResult);

}

カーネル

__global__ void kernel_mean_noshared_3x3(unsigned char* pSrc, unsigned char *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_noshared_3x3(unsigned char* pSrc, unsigned char *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;

}

3×3 平均　シェアード有り

カーネル呼び出し部

//ブロックサイズX
#define BLOCK_SHARED_X 128

void Process_MeanSharedGPU_3x3(unsigned char* pHostSrc, 
unsigned char* pHostResult, int sizex, int sizey, size_t size)
{

	//デバイスメモリのポインタ
	unsigned char* pDevSrc;		//処理前データ
	unsigned char* pDevResult;	//結果

	//ホストメモリをピンロック
	cudaHostRegister(pHostSrc, size, cudaHostRegisterDefault);
	cudaHostRegister(pHostResult, size, cudaHostRegisterDefault);
	//デバイスメモリの確保
	cudaMalloc(&pDevSrc, size);
	cudaMalloc(&pDevResult, size);

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

	//カーネル
	dim3 block(BLOCK_SHARED_X, 1);
	dim3 grid((sizex + BLOCK_SHARED_X - 2 - 1) / ( BLOCK_SHARED_X - 2 ), sizey);

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

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

	//デバイスメモリの開放
	cudaFree(pDevSrc);
	cudaFree(pDevResult);
	//ホストメモリのロック開放
	cudaHostUnregister(pHostSrc);
	cudaHostUnregister(pHostResult);

}

//ブロックサイズX

#define BLOCK_SHARED_X 128

void Process_MeanSharedGPU_3x3(unsigned char* pHostSrc,

unsigned char* pHostResult, int sizex, int sizey, size_t size)

{

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

unsigned char* pDevSrc; //処理前データ

unsigned char* pDevResult; //結果

//ホストメモリをピンロック

cudaHostRegister(pHostSrc, size, cudaHostRegisterDefault);

cudaHostRegister(pHostResult, size, cudaHostRegisterDefault);

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

cudaMalloc(&pDevSrc, size);

cudaMalloc(&pDevResult, size);

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

cudaMemcpy(pDevSrc, pHostSrc, size, cudaMemcpyHostToDevice);

//カーネル

dim3 block(BLOCK_SHARED_X, 1);

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

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

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

cudaMemcpy(pHostResult, pDevResult, size, cudaMemcpyDeviceToHost);

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

cudaFree(pDevSrc);

cudaFree(pDevResult);

//ホストメモリのロック開放

cudaHostUnregister(pHostSrc);

cudaHostUnregister(pHostResult);

}

カーネル


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

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

	unsigned int sum = 0;

	//シェアードメモリ
	__shared__ int BUFF1[BLOCK_SHARED_X]; //上
	__shared__ int BUFF2[BLOCK_SHARED_X]; //中
	__shared__ int BUFF3[BLOCK_SHARED_X]; //下

	//自分の上下３段分をシェアードメモリにセット
	BUFF1[tx] = pSrc[(idy - 1) * sizex + idx];
	BUFF2[tx] = pSrc[(idy)* sizex + idx];
	BUFF3[tx] = pSrc[(idy + 1) * sizex + idx];

	//スレッドを同期
	__syncthreads();

	if (tx < 1 || tx >= BLOCK_SHARED_X - 1) {
		return;
	}

	//シェアードメモリから3x3平均を計算

	for (int x = 0; x < 3; x++) {
		sum += BUFF1[tx - 1 + x];
		sum += BUFF2[tx - 1 + x];
		sum += BUFF3[tx - 1 + x];
	}

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

	return;
}

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

__global__ void kernel_mean_shared_3x3(unsigned char* pSrc, unsigned char *pResult,

int sizex, int sizey)

{

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

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

int tx = threadIdx.x;

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

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

}

unsigned int sum = 0;

//シェアードメモリ

__shared__ int BUFF1[BLOCK_SHARED_X]; //上

__shared__ int BUFF2[BLOCK_SHARED_X]; //中

__shared__ int BUFF3[BLOCK_SHARED_X]; //下

//自分の上下３段分をシェアードメモリにセット

BUFF1[tx] = pSrc[(idy - 1) * sizex + idx];

BUFF2[tx] = pSrc[(idy)* sizex + idx];

BUFF3[tx] = pSrc[(idy + 1) * sizex + idx];

//スレッドを同期

__syncthreads();

if (tx < 1 || tx >= BLOCK_SHARED_X - 1) {

return;

}

//シェアードメモリから3x3平均を計算

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

sum += BUFF1[tx - 1 + x];

sum += BUFF2[tx - 1 + x];

sum += BUFF3[tx - 1 + x];

}

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

return;

}

5×5のコードは、単純に3×3の平均範囲を増やしたもので、長くなってしまうので省略しています。

速度比較

各カーネルの時間は以下の通りです。

	3×3平均	5×5平均
シェアードメモリ使用無し	2.5msec	5.7msec
シェアードメモリ使用有り	2.9msec	4.5msec

5×5でグローバルメモリへのアクセス回数が増えてくると、シェアードメモリを使用したほうが早くなることが判ります。

補足

占有率

シェアードメモリを確保しすぎると、占有率(Occupancy)が下がる可能性があります。

占有率とはWarp(32Threadが1Warp)の最大数と同時実行可能なWarp数の比です。
共有資源を確保しすぎていると、同時実行可能なWarp数は減ります。

高ければ良いわけではありませんが、10・20％など、極端に低い場合はシェアードメモリの使用容量か利用自体を見直したほうが良いでしょう。

占有率はNsightのタイムラインから見ることができます。

シェアードメモリの有効性

以前よりもグローバルメモリへのアクセス速度が上がってきているため、シェアードメモリとの速度差が減ってきています。

シェアードメモリの容量は以前からほとんど増えていない上に、NVIDIAの方のいわく、シェアードメモリを使用しなくても高速にできるようにする方針だそうです。

シェアードメモリを使うとカーネルが複雑化するので、使わないで済むようになってほしいですね。

シェアードメモリの特徴

テスト処理内容

コード例

3×3 平均 シェアード無し

3×3 平均 シェアード有り

速度比較

補足

占有率

シェアードメモリの有効性

3×3 平均　シェアード無し

3×3 平均　シェアード有り