关于是否有重复计算，做无用功的问题

加速计算 CUDA
1

版主你好
我在用CUDA做图像最后的显示的问题，发现这样做最后时间居然比C慢一倍，我猜测是不是出现了线程进行了重复计算，
也就是假如线程2又计算了一遍线程1已经计算完的结果。我的程序结果与C做的相比是正确的。
1.floatdev_sal是将RGB图像变换成灰度图像，进行归一化的输入，
2.floatdev_img3f是最开始的进行归一化的RGB输入图像
3.,int threshold是进行二值化时使用的阈值
4.float*dev_image_output1是最后的输出，宽度是图像的三倍，高不变
5。block是dim3 Dimblock（16,16,1）
grid 是dim3 Dimgrid((width+16-1)/16,(height+16-1)/16,1)

global void Combinekernel(floatdev_sal,floatdev_img3f,int width,int height,int threshold,floatdev_image_output1)
{
int tid_x=threadIdx.x+blockDim.xblockIdx.x;
int tid_y=threadIdx.y+blockDim.y*blockIdx.y;

int tid  =tid_y*width+tid_x;

if((tid_x<width)&&(tid_y<height))
{
       //计算B,G,R的均值
       dev_image_output1[tid_y*3*width+tid_x]=                   
                                                       (dev_img3f[3*tid]+dev_img3f[3*tid+1]+dev_img3f[3*tid+2])*255/3;
      //灰度归一化的乘255
       dev_image_output1[tid_y*3*width+tid_x+width]=dev_sal[tid]*255;
          //对其进行二值化处理
       dev_image_output1[tid_y*3*width+tid_x+2*width]=((dev_sal[tid]*255)>threshold)?255:0;
}

}

2

LZ您好：

看了一下您的代码，您每个线程都只计算一组和自己线程编号相对应的数据，并无重复计算问题。

以及无法判定这个和C的实现是否一致，因为这里只有这一个实现，您需要自己保证。

以及您的运行时间，这个受到多方面的因素影响，如有否受到数据传输影响，比如计算量的影响，以及代码的优化程度等。

大致如此，供您参考。

祝您好运！

3

谢谢版主的回答
我的C段代码测试的时间为1ms，代码如下

  1.float*sal是将RGB图像变换成灰度图像，进行归一化的输入，
  2.float*img3f是最开始的进行归一化的RGB输入图像
  3.,int threshold是进行二值化时使用的阈值
  4.float*image_output1是最后的输出，宽度是图像的三倍，高不变

    clock_t start,finish;

start=clock();

for (int r = 0; r < sal.rows; r++)
{
//为数据的每一行的首地址指针
float *s = sal.ptr(r);
float *s1=image_output1.ptr(r);
float *rgb=img3f.ptr(r);
for (int c = 0; c <sal.cols; c++,rgb+=3)
{
//数据处理部分
s1[c+sal.cols]=s[c]*255;
s1[c]=(rgb[0]+rgb[1]+rgb[2])255/3;
if ((s[c]255)>threshold)
{
s1[c+2sal.cols]=255;
}
else
{
s1[c+2sal.cols]=0;
}
}
}
finish=clock();
double time=finish-start;
printf(“C处理此部分花费的时间所花的时间%fms”,time);

而我将它转成CUDA代码是，内核函数kernel所运行时间0.25ms，但是我测量完整的分配显存拷贝数据和调用kernel的函数的测量总的时间居然达到了67ms，代码如下
/////////////////////////////////////////////////////////////////////////////////

1.floatdev_sal是将RGB图像变换成灰度图像，进行归一化的输入，
2.floatdev_img3f是最开始的进行归一化的RGB输入图像
3.,int threshold是进行二值化时使用的阈值
4.floatdev_image_output1是最后的输出，宽度是图像的三倍，高不变
int Combin(floathos_sal,floathos_img3f,int width,int height,
int threshold,floathos_image_output1)
{
checkCudaErrors(cudaSetDevice(0));

cudaEvent_t start,stop;
float time;
int size=widthheightsizeof(float);

floatdev_sal=NULL;
floatdev_img3f=NULL;
float*dev_image_output1=NULL;

checkCudaErrors(cudaMalloc((void**)&dev_sal,size));
checkCudaErrors(cudaMalloc((void**)&dev_img3f,3size));
checkCudaErrors(cudaMalloc((void**)&dev_image_output1,3size));

checkCudaErrors(cudaMemcpy(dev_sal,hos_sal,size,cudaMemcpyHostToDevice));
checkCudaErrors(cudaMemcpy(dev_img3f,hos_img3f,3*size,cudaMemcpyHostToDevice));

checkCudaErrors(cudaEventCreate(&start));
checkCudaErrors(cudaEventCreate(&stop));
checkCudaErrors(cudaEventRecord(start,0));

dim3 Dimblock(16,16,1);
dim3 Dimgrid((width+Dimblock.x-1)/Dimblock.x,(height+Dimblock.y-1)/Dimblock.y,1);

Combinekernel<<<Dimblock,Dimgrid>>>
(dev_sal,dev_img3f,width,height,threshold,dev_image_output1);
getLastCudaError(“Combinekernel function execution failed”);

checkCudaErrors(cudaEventRecord(stop,0));
checkCudaErrors(cudaEventSynchronize(stop));
checkCudaErrors(cudaEventElapsedTime(&time,start,stop));
checkCudaErrors(cudaMemcpy(hos_image_output1,dev_image_output1,3*size,cudaMemcpyDeviceToHost));

checkCudaErrors(cudaFree(dev_sal));
checkCudaErrors(cudaFree(dev_img3f));
checkCudaErrors(cudaFree(dev_image_output1));

printf(“kernel cost time=%fms\n”,time);
return 0;
}
请问版主，这么做有什么弊端吗，为什么时间会差距的如此之大，希望版主提点宝贵意见

4

(1)第一次的API调用（例如cudaMalloc()), 将花费一定的时间进行cuda runtime初始化。
(2)将数据从内存传输到显存，将需要以PCI-E的速度（甚至更差）进行传输，这也需要大量时间。
(3)然后才是您的kernel运行时间。

根据您的描述，您前2两项将占据66.75ms, 而kernel只有0.25ms

但前2项操作无法避免，我没有更好的建议能让你取消“runtime初始化"以及“显存复制”操作。

如果kernel只运行一次，建议不上CUDA(这样一共才1ms),
而一旦上了CUDA, 前2项固定耗时的操作每个用户都会有，没有办法的。

5

谢谢版主的回答
第一次调用cuda runtime是进行的初始化时间能耗时能达到60多ms啊，我的程序是处理视频的，有1500帧，用C处理单针的时间为43ms，是不是这个初始化的时间只花在处理第一帧时的第一次cuda runtime调用，这样处理视频时能上CUDA吗。

6

嗯嗯。如果您的kernel处理不只一次（一帧），后续就没有开始的初始化和分配空间的开销了。

以后就会很快了（但传输可能无法避免）。

7

我来稍微补充一下横扫斑竹：

1：请LZ注意在最前面申请GPU端的存储空间，在最后释放，中间处理多帧的时候重复使用，而不要将申请和释放打包在某个函数里面，每次都调用。——一般来说LZ不会这样重复申请和释放的，但是以前论坛有过类似的案例，所以再强调一下。

2：多帧的时候，虽然传输时间不可避免，但是LZ可以使用多stream尽量予以掩盖。

大致补充两点注意事项，供LZ参考。

祝您好运~

8

谢谢版主的回答
可是并不是一次性的给视频数据分配空间，也不是将完整的视频数据一次性的由内存拷贝到显存，而是一帧一帧的分配空间，拷贝每一帧，处理完一帧后，还要释放这部分空间，您所说的后续部分没有分配空间的开销了，这部分有点不理解

9

为何每帧都重复的分配-释放呢？

既然上一帧处理完了，它的空间就可以重用了，没必要立刻cudaFree+cudaMalloc的。

您觉得呢？

当然，如果贵实验室、单位有此方面的编码习惯要求（即每个缓冲区不得重用，要立刻释放+重新分配）则上文建议请无视。

10

ice版主和横扫版主的话让我茅塞顿开啊，
也就是说（帧的循）环处理前处理之前就把显存的空间分配好，不要再做帧循环处理时分配空间，还有在所有循环结束后再释放显存空间，期间这部分空间一直可以被每一帧使用。但是拷贝数据必须是每一帧每一帧的拷贝吧

11

谢谢版主的回答

12

LZ您好：

1：是的，尽量复用即可，但是每次必须copy数据进去。

2：以及您可以使用多个stream互相掩盖计算和传输，此时需要申请多个存储空间，你使用几个stream就申请几个空间，以及host端需要使用页锁定内存。

大致如此，祝您编码顺利~

13

明白了，谢谢版主的回答

14

不客气的，我和横扫版主将竭诚为您服务。

祝您编码顺利~