[table=98%]
[tr][td]函数需要完成对矩阵的每个点对其r的行领域求和，边界做镜像对称，即若第一行数为1 2 3 4 5 6 7，r=2，则结果为 11 12 15 20 25 28 29，现要对WH的矩阵进行此操作，分配H个块，在块内分配W个线程，即每个点的计算对应一个线程，每个块对应每行的计算，在kernel函数里面动态分配共享存储器extern shared float si[]，
由于涉及到边界的处理，因此分配块内si的大小为W+2r,即si里面存放这一行点的数据，利用si对其进行行领域求和操作，外部调用核函数为kernel<<<H, W, H*(W+2*r)*sizeof(float)>>>(**);
但是程序没有得到正确的结果，不知道是不是共享存储器的使用出现了问题？[/td][/tr]
[/table]
LZ您好：
仅凭您上述叙述，我无法确定问题所在，甚至无法理解您的实现意图，所以无法给出您建议了。
请您重新叙述并给出代码，这样才能对您的问题加以判断。
以及，使用shared memory请注意使用__syncthreads()函数，以保证可靠写入后再分享数据。
大致如此，祝您好运~
/对矩阵进行行均值滤波，即求出每个像素点关于其(2r+1)行领域像素的和值
由于每行的像素点个数W可能超过块内最大线程数，因此将W分为n个，
线程数分为W/n，每个线程处理一个像素点的计算，一个块里处理W/n个像素点，共享存储器si存储
的数据为前面r个存放首部扩展的边界数据，中间的W/n个存放对应的像素点数据，后面的r个存放尾部
扩展的边界数据，即每个块里的si存放(W/n+2*r)个数据
*/
global void
d_meanfilter_x_shared(float *id, float od, int w, int r)
{
extern shared float si[];
unsigned int tid = threadIdx.x;
unsigned int gid = threadIdx.x + blockIdx.xblockDim.x; //对应矩阵下标索引
float sum=0;
si[tid+r] = id[gid];
if(tid==0 && gid%w==0) //矩阵每行的第一个像素且为块内第一个线程
{
for(int i=0; i<r; i++)
si[r-1-i] = id[gid+i+1];
}
else if(tid==0 && gid%w!=0)
{
for(int i=0; i<r; i++)
si[i] = id[gid-r+i];
}
else if(tid==(blockDim.x-1) && (gid+1)%w!=0)
{
for(int i=0; i<r; i++)
si[tid+r+i+1] = id[gid+1+i];
}
else if(tid==(blockDim.x-1) && (gid+1)%w==0) //矩阵每行的最后一个像素且为块内最后一个线程
{
for(int i=0; i<r; i++)
si[tid+r+i+1] = id[gid-1-i];
}
__syncthreads();
for(int i=0; i<=2*r; i++)
{
sum += si[tid+i];
}
__syncthreads();
od[gid] = sum;
}
核函数调用为：
d_meanfilter_x_shared<<<bmpHeight2, bmpWidth/2, bmpHeight2*(bmpWidth/2+2*radius)*sizeof(float)>>>(d_Y, d_temp, bmpWidth, radius);
LZ您好：
目测了您3#给出的代码，似乎并无问题。
以及，您4#调用参数可能有点问题：“bmpHeight2(bmpWidth/2+2*radius)sizeof(float)”，这个是一个block的shared memory的使用量，考虑到您一个block只有“bmpWidth/2”个线程，那么“(bmpWidth/2+2radius)*sizeof(float)”这个数量应该是合适的每个block的shared memory的申请数量。
以及，shared memory的数量是有限的，申请过多可能导致kernel无法运行。
最后，这里假定您的bmp的宽度的数值和合适的，并没有这里的整数除法而少算一段数据。
大致如此，供您参考。
您还可以在kernel后面调用cudaDeviceSynchronize()函数，并检查该函数的返回值，看看究竟提示是什么问题。
祝您编码顺利~
确实是参数调用的问题，换成一个block的shared memory的使用量才是正确的，但是速度却变慢了，原来的做法是将数据放在全局存储器里，一个线程对应一个点，个人感觉程序变慢的原因应该是核函数里面有if/else判断造成，而实际的共享存储器肯定比全局存储器访问速度快，被这个问题困扰了好长时间，非常感谢版大的热心解答！
再次感谢！！！
LZ您好：
恭喜您解决了原程序的问题。
如果您继续打算使用shared memory的话，不妨如此试下：
1：一个block使用W/n+2*r个线程一次性读入数据到shared memory内部并同步。
2：一个if使前r个线程和最后r个线程return，中间的W/n个线程干活并输出。
这样或许可以简化逻辑，达到更好的效果。
祝您编码顺利~
斑竹您好：
这个貌似在读入数据到shared memory的时候也涉及到判断呀，因为r里面存放的是边界或者重叠数据，所以还是得判断；
我想问一下版大，核函数里面涉及到if/else判断的是否会严重影响程序的速度？如果没有判断的话，这里使用共享存储器是否一定比全局存储器快呢？
我替ICE回答一下，ICE可能中暑了…
if…else…是否会影响效率主要看2点：
（1）您的if的条件多长？如果是if (a > b || c > d || e > f… || y > z)这么一串的话，而if和else的body又不长，那果断还是会影响一点的。
（2）当然，1种夸张了，说下一般情况，如果您的if…else能大约分界在warp的边界，那么基本不会影响什么。否则，如果分支在warp内部，可能会造成严重影响。
感谢来访，这么热的天，还是刮掉胡子吧。
LZ您好：
关于if-else判断的开销问题，请参考9#横扫千军斑竹的详细说明。
以及，你看一下你if判断的那四个部分，每部分都是一个线程干活，循环赋值。我的意思是说，你把这个变成一组线程，也就是把这个循环用线程展开，应该比一个线程循环效果好一些，特别是如果r稍大一些的时候。
大致如此，祝您编码顺利~
真是涨知识鸟！！谢谢版大呦！！PS:就喜欢大胡子的沧桑赶脚！
嗯嗯，我下去实现一下这种方法！
非常感谢您！