图像的均值滤波中用到共享存储器出现了问题

[wkdxiad2008]

假如图像高度为H宽度为W，均值滤波思路为先对行进行滤波再对列进行滤波，对行滤波时，考虑每行W个像素点开W个线程形成一个块，分配2*W个动态共享存储器单元，前W个存放像素点原始值，后W个存放对应像素点滤波之后的值，整幅图有H个块，滤波边界做“replicate”填充，对应代码如下：
__global__ void
d_mfilt_x_shared(float *id, float *od, int w, int h, int r)
{
extern __shared__ float si[];
unsigned int tid = threadIdx.x;
unsigned int gid = threadIdx.x + blockIdx.x*blockDim.x;

//将数据从global读入shared memory，读入数据后进行一次同步，保证计算时所有数据均已到位
si[tid] = id[gid];
__syncthreads();
if(tid <= r)
{
float t = si[0]*(r+1-tid);
for(int i=1; i<=tid+r; i++)
t += si[i];
si[tid+blockDim.x] = t;
}
else if(tid < w-r)
{
float t = 0;
for(int i=tid-r; i<=tid+r; i++)
t += si[i];
si[tid+blockDim.x] = t;
}
else
{
float t = si[tid+w-1-tid]*(r+tid+2-w);
//float t = 0;
for(int i=tid-r; i<w-1; i++)
t += si[i];
si[tid+blockDim.x] = t;
}
__syncthreads();
od[gid] = si[tid+blockDim.x];
}
这里只是对像素点的行窗口领域求取和值，得到的结果却很怪异，当W=H时能得到正确结果即线程块数和块内线程数相等，但W!=H时就得不到正确的结果了，调了蛮长时间也不知道什么地方出错了，猜测应该是发生了bank conflict，不知道怎么解决这个问题，求大神知道一二，万分感谢

[wkdxiad2008]

本人刚学习CUDA不久，纯新手，大神们多多给点建议啊！！

[linxxx3]

bank conflict 不会影响结果的正确性，应该是程序写错了，不要乱猜么…… 这里拿出来的代码都跟H没有一点关系，最可能出错在外面，不介意的话，贴一下外面调用这个kernel的代码。 另外，保证用来比对的结果，和比对的过程没有问题。

[wkdxiad2008]

嗯 确实跟H没有关系，程序调试出来了，核函数没有问题，外面的程序出了点小问题，谢谢楼上热心的解答！

[wkdxiad2008]

引用 2 楼 linxxx3 的回复:

bank conflict 不会影响结果的正确性，应该是程序写错了，不要乱猜么…… 这里拿出来的代码都跟H没有一点关系，最可能出错在外面，不介意的话，贴一下外面调用这个kernel的代码。 另外，保证用来比对的结果，和比对的过程没有问题。

这种方法的缺陷是需要在一个块里面开W个线程，当W>512时程序得不到正确的结果，但是我的显卡配置为最大块内线程数为1024，W=576时程序错误，W=400时程序正确。要用到共享存储器的话好像回避不了这个问题，真不知道怎么解决。