shared memory使用

加速计算 CUDA
1

我对一幅图像做了2维FFT变换后，得到复数，想对其求模值，分别放在global memory和shared memory中进行计算，但结果是global memory的速度要比shared memory的速度快，不知道是不是我用shared memory方法不对，请高手指教，谢谢！！！

图像大小8192*2048

global memory的代码：
sqrtCaculate<<<512,256>>>(d_signal,d_p,81922048);
static global void sqrtCaculate( Complex a, int* p,int size)
{
const int numThreads=blockDim.xgridDim.x;
const int threadID=blockIdx.xblockDim.x+threadIdx.x;
for(int i=threadID;i<size;i+=numThreads)
{
p[i]=ComplexSqrt(a[i]);
if(p[i]>255)
p[i]=255;

}
}

static device host inline int ComplexSqrt(Complex b)
{
int c;
c=(int)sqrt(b.xb.x+b.yb.y);
c=c/2000;
return c;
}

shared memory代码：

BLOCK_DIM=16;
size_w=8192;
size_h=2048
dim3 grid(size_w/BLOCK_DIM,size_h/BLOCK_DIM,1)
dim3 threads(BLOCK_DIM,BLOCK_DIM,1);
static global void sqrtCaculate( Complex* a, int* p,unsigned int size_w,unsigned int size_h)
{
shared Complex block[BLOCK_DIM][BLOCK_DIM+1];
unsigned int xIndex=blockIdx.xBLOCK_DIM+threadIdx.x;
unsigned int yIndex=blockIdx.yBLOCK_DIM+threadIdx.y;
unsigned threadID=yIndex*size_w+xIndex;
if ((xIndex<8192)&&(yIndex<2048))
{

block[threadIdx.x][threadIdx.y]=a[threadID];
}
__syncthreads();

if ((xIndex<8192)&&(yIndex<2048))
{
p[threadID]=ComplexSqrt(block[threadIdx.x][threadIdx.y]);
}
__syncthreads();

if (p[threadID]>255)
{
p[threadID]=255;
}
}
static device host inline int ComplexSqrt(Complex b)
{
int c;
c=(int)sqrt(b.xb.x+b.yb.y);
c=c/2000;
return c;
}

2

1、对a的数据访问只有一次，没必要用shared memory
相当于：1）读取a；2）读取a到shared memory，读取shared memory
1）更快些
2、block[threadIdx.x][threadIdx.y]=a[threadID]
这个会有bank conflict，应该是block[threadIdx.y][threadIdx.x]=a[threadID]

另
block[BLOCK_DIM][BLOCK_DIM+1]，没注意到你这有个+1，这样就没有bank conflict了
假设block[16][16]的话，shared memory存储有16个bank，行x，列y，数据是这样存储的
bank对应y：0 1 2 …… 15
x=0：0 1 2 …… 15
x=1：16 17 18 …… 31
……
x=15 240 241 242 …… 255
你这样访问：block[threadIdx.x][threadIdx.y]=a[threadID]，假设0-15号线程访问的是block[0][y]——block[15][y]（上面的数据0、16、……240），都是访问的不同行，同一列，也就是bank N位置上的数据，就有bank conflict了。
如果block[16][17]的话：
bank对应y：0 1 2 …… 15
x=0：0 1 2 …… 15
x=1：null 16 17 18 ……
31 null 32 ……
……
x=15 null 240
241 242 …… 255
因为有17列，多了一列数据null，所以16个thread访问的数据（0、16、……240）处于不同bank了，就没有bank conflict了

3

1、对a的数据访问只有一次，没必要用shared memory
相当于：1）读取a；2）读取a到shared memory，读取shared memory
1）更快些
2、block[threadIdx.x][threadIdx.y]=a[threadID]
这个会有bank conflict，应该是block[threadIdx.y][threadIdx.x]=a[threadID]

另
block[BLOCK_DIM][BLOCK_DIM+1]，没注意到你这有个+1，这样就没有bank conflict了
假设block[16][16]的话，shared memory存储有16个bank，行x，列y，数据是这样存储的
bank对应y：0 1 2 …… 15
x=0：0 1 2 …… 15
x=1：16 17 18 …… 31
……
x=15 240 241 242 …… 255
你这样访问：block[threadIdx.x][threadIdx.y]=a[threadID]，假设0-15号线程访问的是block[0][y]——block[15][y]（上面的数据0、16、……240），都是访问的不同行，同一列，也就是bank N位置上的数据，就有bank conflict了。
如果block[16][17]的话：
bank对应y：0 1 2 …… 15
x=0：0 1 2 …… 15
x=1：null 16 17 18 ……
31 null 32 ……
……
x=15 null 240
241 242 …… 255
因为有17列，多了一列数据null，所以16个thread访问的数据（0、16、……240）处于不同bank了，就没有bank conflict了