shared memory使用

我对一幅图像做了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.x
blockDim.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.y
BLOCK_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;
}

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了

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了