09/12/06 01:25:46
CUDA勉強中の者ですが共有メモリの利用で躓いてるところです。アドバイス頂けたら幸いです。
下記の二つのカーネルでCUDA_karnel_sの方が5倍時間がかかってしまうのですが原因がわかりません。
違いは読み込んだデータをグローバルメモリに保存するか共有メモリに保存するかです。
__global__ void CUDA_karnel_g(uchar4 *vram, int sx, int sy, uint1 *vram2)
{
int i = threadIdx.x;
vram += blockIdx.y*sx;
vram2[i] = ((uint1 *)vram)[i];
__syncthreads();
uchar4 px;
*((uint1 *)&px) = vram2[i];
unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x);
px.z = px.y = px.x = Y;
vram[i] = px;
}
__global__ void CUDA_karnel_s(uchar4 *vram, int sx, int sy, uint1 *vram2)
{
int i = threadIdx.x;
vram += blockIdx.y*sx;
__shared__ uint1 shared[125*32];
shared[i] = ((uint1 *)vram)[i];
__syncthreads();
uchar4 px;
*((uint1 *)&px) = shared[i];
unsigned char Y = (unsigned char)(0.2126f*px.z+0.7152f*px.y+0.0722f*px.x);
px.z = px.y = px.x = Y;
vram[i] = px;
}
まずバンクコンフリクトを疑ったのですがCUDA Visual Profilerでみるとwarp serializeは0で発生していませんでした。
意図的にバンクコンフリクトを発生させると更に１０％程遅くなるのでバンクコンフリクトは原因ではなさそうです。
共有メモリは速いはずなのにグローバルメモリよりなぜ遅くなるのか悩んでいます。初歩的なミスだろうとは思うのですが。