关于kernel函数中printf()执行顺序的问题

加速计算 CUDA
1

自己写了一段小代码，因为看好多资料都说GPU的一个thread相对较弱，所以想知道一下它到底有多弱，当然我的这个例子也不能说明什么问题，一个是GPU的运算时间包含了数据传输的时间，另一个原因是我的累加算法实在是太小了，恐怕对于GPU的一个线程来说也没有太大的负担，但是在运行的过程中，我却发现了另一个问题：
为什么kernel函数中的执行顺序不是按照我理解的来的呢？代码如下：


 #include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <time.h>
float cpu_sum(int N)
{
 float sum=0.0;
 for(int i=0;i<N;i++)
 {
  sum=sum+i;
 }
 printf("CPU的计算结果：%f\n",sum);
 return sum;
}
__global__ void gpu_sum(int* M)
{
 float gsum=0.0;
 for(int i=0;i<M[0];i++)
 {
  gsum=gsum+i;
 }
 printf("gpu result is:%f\n",gsum);
}

void main()
{
 clock_t t0,t1,t2,t3;
 t0=clock();
 float csum=cpu_sum(10000);
 t1=clock();
 printf("%d\n",t1-t0);
 int h[1];
 h[0]=10000;
 int *dev;
 t2=clock();
 cudaMalloc((void**)&dev,sizeof(int));
 cudaMemcpy(dev,h,sizeof(int),cudaMemcpyHostToDevice);
 gpu_sum<<<1,1>>>(dev);
 t3=clock();
 printf("%d\n",t3-t2);
 cudaFree(dev);
}

执行的结果如图：
[attach]3006[/attach]
按照我的理解，它不是应该先执行kernel中的print（），然后再执行 printf(“%d\n”,t3-t2);吗？难道是异步执行了？
我设想的顺序应该是CPU的计算结果。。。、0、gpu result is。。。、65，最后两个的显示顺序和我想的相反！

2

顺便问一下，我把整个程序都加载到GPU上运行，包括串行的函数，因为有时候可以避免数据的反复传输，是不是就相当于我在用GPU的一个线程处理整个串行部分呢？这样会不会达不到想要的效果呀！

3

LZ您好，因为device上的printf不是也无法直接控制host端的输出，他是通过某种后台同步机制将数据传回来显示的。

因为该同步机制会有延迟，所以会出现您的情况，以及有时甚至会出现不显示的情况。如果遇到不显示的情况，您可以添加cudaDeviceSynchronize()或者cudaDeviceReset()等同步函数。

类似的问题以前也有讨论，如您有兴趣，不妨翻翻旧帖。

祝您编码顺利~

4

1：推荐尽量把工作都放在GPU上执行，避免来回copy数据。
2：把工作都放在GPU上执行也可以是CPU连续发射kernel，并非一定要GPU程序自己反复循环。

3：无法理解您“相当于XXXXX”的含义，也无法评价“是否达到预期效果”。

大致如上，供您参考。

5

好的，可能是我的理解有误。O(∩_∩)O谢谢！

6

我补充一下4#的第三点

这一点究竟如何，其实是和您具体的实现有关的，所以不能凭空下结论，是否“相当于”，是否“效果好”。要具体问题具体分析才行。

大致如上，供您参考。

7

虽然CUDA后期加入了printf功能，但是感觉很鸡肋。一般都是建议不要再GPU端用printf。因为printf的执行顺序很没有保障。你的程序实质上只有一个gpu线程，要是多线程，printf会让你抓狂的。正常检验GPU计算是否正确都是先将GPU数据传回CPU再输出比较的。

8

感谢您参与讨论！

我觉得一般来说，可以copy回来检验，也可以nsight调试查看，对于一些简单的反馈结果kernel里面printf是比较直观和简单的，也可以使用。

9

kernel会在开始执行之后立即将控制交给CPU线程执行后面函数。你的例子中，CPU printf执行在kernel结束之前所以会先显示CPU的printf。