我的程序有两个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差了一倍。
要么楼主此图是假的，要么楼主前文说慌了。
果断不是一样的参数编译的。
来问问题，我给你痛快解答，你也不能消遣我们啊。
[
是同样的代码，没改过，就移动了位置，肯定是编译有不同嘛，问题就是为什么会编译不一样，因为我其他都没改过。