我的程序有两个cu文件,为啥我有一个kernel函数放在一个cu文件里时间为5ms,在另一个cu文件里时间为20ms呢?同一个函数,代码完全一样。
尚不清楚您这个情况是何原因,我觉得执行时间应该没有差别才对。
我猜测可能是计时上的问题。
欢迎其他人补充,以及LZ继续提供进一步的信息。
预祝您春节愉快~
[
有两个cu文件 a.cu 和a_kernel.cu
其中cu端包含有
函数里面调用
extern “C” void cuda_stereomatching(uchar *imgL,uchar *imgR,float *supdisp,uchar *dispvalid)
{
…
crossconstruct<<<blockH,threadW>>>(imgL,dev_armL ); //要测试的kernel函数
…
}
代码如下:
#define max3(x,y,z) max(max(x,y),z)
#define absdiff(x,y) ((x)>(y)?(x)-(y): (y)-(x))
global void crossconstruct(uchar imgdat,uchar arm) //480 640
{
shared uchar imgrow[MAX_W3];
short x=threadIdx.x;
short y=blockIdx.x;
int tid=yblockDim.x+x;
int i;
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[3x+1];
uchar red =imgrow[3x+2];
uchar aa = 0;
for(i=1;i<L1&&i<=x;i++)
{
aa = max3(absdiff(blue,imgrow[3(x-i)]),absdiff(green,imgrow[3*(x-i)+1]),absdiff(red,imgrow[3*(x-i)+2]));
if(aa>=(i<L2?Tao1:Tao2))
{
break;
}
}
arm[tid]=i-1;
for(i=1;i<L1&&i<IMG_W-x;i++)
{
aa = max3(absdiff(blue,imgrow[3*(x+i)]),absdiff(green,imgrow[3*(x+i)+1]),absdiff(red,imgrow[3*(x+i)+2]));
if(aa>=(i<L2?Tao1:Tao2))
{
break;
}
}
arm[tid+IMG_H*IMG_W]=i-1;
for(i=1;i<L1&&i<=y;i++)
{
aa = max3(absdiff(blue,imgdat[3*(tid-iIMG_W)]),absdiff(green,imgdat[3(tid-iIMG_W)+1]),absdiff(red,imgdat[3(tid-iIMG_W)+2]));
if(aa>=(i<L2?Tao1:Tao2))
{
break;;
}
}
arm[tid+2IMG_HIMG_W]=i-1;
for(i=1;i<L1&&i<IMG_H-y;i++)
{
aa = max3(absdiff(blue,imgdat[3(tid+iIMG_W)]),absdiff(green,imgdat[3(tid+iIMG_W)+1]),absdiff(red,imgdat[3(tid+iIMG_W)+2]));
if(aa>=(i<L2?Tao1:Tao2))
{
break;
}
}
arm[tid+3IMG_H*IMG_W]=i-1;
}
将crossconstruct函数放a.cu和a_kernel.cu里面效率居然不一样,其它函数都一样的,就这个不一样,很奇怪。计时没问题的。
将__global__放在某个特定的.cu和其他.cu中,在同样的编译参数下,不会导致楼主的一个5ms即可,放在另一个.cu需要20ms才行的现象。甚至你单独将cubin提取出来,放在单独的另外一个磁盘文件中,也不会影响该cubin里的kernel的实际运行速度的。
但是,楼主的确说你观测到了这种现象,这是为何?可能的原因有多个:
(1)可能是您对时间的测试上的问题。(如同ICE大所说)
(2)2个文件的编译参数不同。
一般会是这2种情况,当然,也不排除楼主(3)眼花了的因素。
还请主要检查(1)(2)情况。
预祝春节愉快!
时间测试上没有问题的
参数是一样的
cuda_starttime(“aggregateH”);
crossconstruct<<<blockH,threadW>>>(imgL,dev_armL ); //20ms
cuda_endtime();
我这个调用时放在a.cu里面
然后在a.cu里声明
global void crossconstruct(uchar *imgdat,uchar *arm);
在a.cu
里面定义函数
global void crossconstruct(uchar *imgdat,uchar *arm)
{
…
}
其过程执行时间为5ms
然后
a.cu里面声明改为
extern global void crossconstruct(uchar *imgdat,uchar *arm);
把a.cu里面的
crossconstruct(uchar *imgdat,uchar *arm);注释掉,复制到a_kernel.cu 里面时间就变为20ms了。其他都没变,百思不得其解。
我把计时函数也贴出来好了:
void cuda_starttime(char *str)
{
strcpy(cuname,str);
if(initflag==0)
{
cudaEventCreate(&start);
cudaEventCreate(&stop);
initflag=1;
}
cudaEventRecord(start,0);
}
void cuda_endtime()
{
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&etime,start,stop);
cout<<“(”<<cuname<<“)”;for(int i=0;i<20-strlen(cuname);i++) cout<<" ";cout<<“elapsed time:”<<etime<<“ms”<<endl;
}
两次的运行结果是一样的,就显示时间不一样。
请问楼主,是对kernel的首次运行进行时间测试的吗?
如果是,请放弃初次结果,而是用第二次的结果。
第一次的极度不准。原因众多。你可以测试第二次。
楼主的kernel调用后有加cudaThreadSynchronize()吗?貌似没看到~
这个侧过好多次的,都是的
这个在主机函数里面有加的,在CPP文件里的main函数加了
嗯嗯。那就奇怪了。不知道为何忽然会减慢4x了。
那个不需要的。他对流0里的event进行了同步,等效于device同步。建议tianyuan下次仔细看。
哦,对哦,他的计时函数里面有同步!
就是,我就很奇怪。其他kernel函数都不会,就这个kernel产生这种情况。纳闷了
时间是没有错的,我用NVPP分析过时间了
嗯嗯。对楼主的建议,都是在假设楼主描述完全无误的情况下进行的。
既然在visual profiler里看到kernel的确一次运行5ms, 在移动位置重新编译后却需要运行20ms。那么最大的可能就是这2次编译出的代码不同,(因为如果编译出的代码一样,同样的参数不可能突然变慢400%的),而编译出的代码不同,那么最大的可能就是编译参数的问题。
但是楼主已经在前文否定了此情况,那么在我的理解层次内,我无法给出其他建议了。
建议原厂支持介入。
[attach]2927[/attach][attach]2926[/attach]
这是用NVVP分析的结果,显示两者运行不一样,而其他的kernel函数运行分析到的是一样
这个真的是同样的代码,同样的参数么?
Branch Divergence Overhead差了一倍。
要么楼主此图是假的,要么楼主前文说慌了。
果断不是一样的参数编译的。
来问问题,我给你痛快解答,你也不能消遣我们啊。
[
是同样的代码,没改过,就移动了位置,肯定是编译有不同嘛,问题就是为什么会编译不一样,因为我其他都没改过。