共享存储器的问题

[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+2
r,即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.x
blockDim.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:就喜欢大胡子的沧桑赶脚!

嗯嗯,我下去实现一下这种方法!

非常感谢您!