shared memory使用

[小鱼的dream]

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

图像大小8192*2048

global memory的代码：
sqrtCaculate<<<512,256>>>(d_signal,d_p,8192*2048);
static __global__ void sqrtCaculate( Complex* a, int* p,int size)
{
const int numThreads=blockDim.x*gridDim.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.x*b.x+b.y*b.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.x*BLOCK_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.x*b.x+b.y*b.y);
c=c/2000;
return c;
}

[zjx1020]

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]

[小鱼的dream]

[Quote=引用 1 楼 zjx1020 的回复:]
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]
[/Quote]
我对bank conflict不是怎么太懂，看书也没怎么看明白，我的QQ:764789211,希望能多指教，呵呵

[zjx1020]

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了

[小鱼的dream]

[Quote=引用 3 楼 zjx1020 的回复:]
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 ……
[/Quote]
谢谢呀，能加你qq吗？想多向你请教，呵呵！我的qq:764789211