kernel 放在不同的cu文件会影响效率吗?

我的程序有两个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_W
3];
short x=threadIdx.x;
short y=blockIdx.x;
int tid=y
blockDim.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[3
x+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+2
IMG_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+3
IMG_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差了一倍。

要么楼主此图是假的,要么楼主前文说慌了。

果断不是一样的参数编译的。

来问问题,我给你痛快解答,你也不能消遣我们啊。

[

是同样的代码,没改过,就移动了位置,肯定是编译有不同嘛,问题就是为什么会编译不一样,因为我其他都没改过。