高手帮看看这个代码

global_ void crossconstruct(uchar imgdat,uchar arm) //480 640
{
shared uchar imgrow[MAX_W
3];
shared uchar imgrowpre[MAX_W
3];
shared uchar imgrowlast[MAX_W3];
//
short x=threadIdx.x;
short y=blockIdx.x;
int tid=y
blockDim.x+x;
uchar i;
uchar l=0,r=0,u=0,d=0;
uchar aa[4];

imgrow =imgdat[3yIMG_W+x];
imgrow[IMG_W+x] =imgdat[3yIMG_W+IMG_W+x];
imgrow[2IMG_W+x]=imgdat[3yIMG_W+2IMG_W+x];
__syncthreads();
uchar blue =imgrow[3x];
uchar green=imgrow[3
x+1];
uchar red =imgrow[3*x+2];
for(i=1;i<L1&&i<=x;i++)
{
if(colordis(imgrow,blue,green,red,x-i)>=(i<L2?Tao1:Tao2))
{
break;;
}
}
aa[0] = i-1;

for(i=1;i<L1&&i<IMG_W-x;i++)
{
if(colordis(imgrow,blue,green,red,x+i)>=(i<L2?Tao1:Tao2))
{
break;
}
}
aa[1] = i-1;

for(i=1;i<L1&&i<=y;i++)
{
if(colordis(imgdat,blue,green,red,tid-i*IMG_W)>=(i<L2?Tao1:Tao2))
{
break;;
}
}
aa[2] = i-1;

for(i=1;i<L1&&i<IMG_H-y;i++)
{
if(colordis(imgdat,blue,green,red,tid+i*IMG_W)>=(i<L2?Tao1:Tao2))
{
break;
}
}
aa[3] = i-1;

arm[tid]=aa[0];
arm[tid+IMG_HIMG_W]=aa[1];
arm[tid+2
IMG_HIMG_W]=aa[2];
arm[tid+3
IMG_H*IMG_W]=aa[3];

}
不要后面四个
arm[tid]=aa[0];
arm[tid+IMG_HIMG_W]=aa[1];
arm[tid+2
IMG_HIMG_W]=aa[2];
arm[tid+3
IMG_H*IMG_W]=aa[3];
整个程序效率大概是0.05ms左右,加上了就20ms左右
crossconstruct<<<blockH,threadW>>>(imgL,dev_armL ); BLOCKH为480 threadw为640

cudaMalloc((void **)&dev_armL,HW4);
就算是全局变量效率也没那么低吧

arm[tid]=aa[0];
arm[tid+IMG_HIMG_W]=aa[0];
arm[tid+2
IMG_HIMG_W]=aa[0];
arm[tid+3
IMG_H*IMG_W]=aa[0];
换为这样大概是6ms

换为这样:
arm[tid]=1;
arm[tid+IMG_HIMG_W]=1;
arm[tid+2
IMG_HIMG_W]=1;
arm[tid+3
IMG_H*IMG_W]=1;
0.079ms
快要疯了,求救

大致看了一下您的代码,您在kernel里面的uchar aa[4]数组应该是放在local memory里面的,详细的瓶颈分析建议跑一下nvvp(nvidia visual profiler)看看。

不过如果寄存器数量没什么压力的话,您可以将uchar aa[4]改写为 uchar aa0,aa1,aa2,aa3;这样就放在寄存器里面了,您不妨先试验一下。

祝您编码愉快。

过如果寄存器数量没什么压力的话,您可以将uchar aa[4]改写为 uchar aa0,aa1,aa2,aa3;这样就放在寄存器里面了,您不妨先试验一下。
这个改了,还是一样的20ms

那就无法直接目测得出结果了,您不妨跑跑nvvp看看是怎么一个性能分析。

好的,我试试,不过没用过这个工具

嗯嗯,推荐试一下的。

有没有帖子介绍怎么用的?

要下载安装么?还是装CUDA 5.0自带了的

应该是CUDA Toolkit自带的,您可以先参考一下自带的手册看看。

[

如图,只有file这一项是required的,图中已经显示为需要给出一个可执行文件,这个就是你CUDA工程编译得到的exe文件。

祝您编码愉快~

我运行完了,这个函数分析
你看看瓶颈在哪?
要看哪几项的?

问题就在图中两个感叹号的地方,一个是global的读取效率低下,一个是分支开销。

前者需要您尽量解决合并访问的问题,后者您可以看看是否可以把分支做成warp对齐的。

祝您编码愉快~

好的,谢谢啦。我再看看。

合并访问是不是一般用纹理,或者多位合成32位来处理的?
我把4为uchar合成一位int处理速度好像也是一样啊

合并访问不是这个意思,详情请参阅一下programming guide等官方资料。

这个该如何实现合并访问呢?能给个大概思路吗?我看了挺多资料,但没看有什么例子,感觉挺模糊的。

您可以看一下programming guide的附录F,这里面有针对不同架构下,global memory访存的特点。

在fermi和kepler这样带有L2 cache的GPU下,对合并访问的限制已经大为放宽了,一般只要保证连续的线程访问较为集中分布的数据即可,无需要求顺序性,也基本无需考虑对齐,即可达到较好的效果。

大致如上,祝您编码顺利。