关于GPU与CPU并行的问题

由于GPU的内核执行以及一些数据传输是异步的,故为GPU和CPU之间的并行提供了可能性。但貌似仍然有个问题:
如果我理解没错的话,CPU的代码必须写在GPU内核/异步数据传输和GPU其他具有阻塞性质的API之间才能实现CPU/GPU并行。
如:
1:cudaMalloc();
2:cudaMemcpyAsyn();
3:kernel<<<>>>();
4:cudaMemcpyAsyn();
5:CPUfoo();
6:cudaFree();
然而如果将5、6句互换位置,CPUfoo应该就不能和GPU端并行了吧?

当然这可能对于自己维护的cuda代码不是什么大问题。
但是,现在的GPU程序大多数还是从原来的CPU程序中改过来的,为了与原程序连接,要保持原程序接口,故GPU代码不得不封装成如下形式:
GPUfoo( cpu parameter){
cudaMalloc(gpu parameter);
cudaMemcpyAsyn(cpu->gpu);
kernel<<<>>>(gpu);
cudaMemcpyAsyn(gpu->cpu);
cudaFree(gpu parameter);
return status;
}
这样的话,GPUfoo与CPUfoo是不是就不能并行了?

当然,我也考虑过应用CPU端多线程+共享状态变量的形式,一个线程调用GPUfoo,一个调用CPUfoo的方法,但总觉得太曲线救国了吧?

在讨论您的所有的2个例子之前,我先说一下,CPU和GPU并行有2点。

其中一点是对于特殊的负责和GPU交互的线程来说,此时你就需要小心了(例如你的上文的尝试)。
而另外一点则是传统的和GPU无关系的host thread, 这时候它该继续和其他host thread并行,那就会继续并行。

所以,最简单的并行方式就是:
多线程(host上), 一个/多个线程在CPU上计算,压榨CPU;
一个/多个线程在GPU行算,不用考虑CPU代码能否并行,能否充分利用CPU(由前者负责压榨,这专心交流GPU即可)。

不过很多人不想启动额外的线程,想利用异步API, 见缝插针的利用上CPU. 也是可以的,这估计也是楼主喜欢的,我继续说。

------------------有限长度的分割线-------------------

首先说,楼主的代码1里面的代码是一般情况下无问题的,因为cudaFree会负责自动同步默认流一次,然后会导致和你的cudaMemcpyAsync所在的流也自动同步一次。
(为了您的安全,例如您使用了不和默认流同步的流的时候,我建议您在cudaMemcpyAsync后手工同步一次,这里就不继续讨论了,因为不是您的主题)

在我们达成“一般情况下该代码无问题”的时候,的确CPUfoo();和cudaFree()颠倒就不能见缝插针的利用了。因为此时会立刻等待。所以您的假设是合理的。

------------------有限长度的分割线-------------------
然后继续看您的代码2,是的,如果您直接
GPUfoo();
CPUfoo();
那么此时他们就不能同步执行了,
因为您的GPUfoo里面是一个完整的执行-等待流程,当CPUfoo()开始之前后,实际上前者已经执行完毕了。所以您的假设依然正确。

那么如何才能实现楼主的目标了,楼主您给出了一个“曲线救国”的方式,不过在我看来,这是正规方式。现在的年轻人写代码都喜欢这么写。

以及,您依然可以用传统的方式,不用新开host线程, 您可以尝试进行回调(call back).
void GPUfoo(void (*cpu_callback)(void *), void *work_data)
{
cudaMalloc(gpu parameter);
cudaMemcpyAsyn(cpu->gpu);
kernel<<<>>>(gpu);
cudaMemcpyAsyn(gpu->cpu);

//现在命令已经发布完毕了,可以等待了,那么在此回调CPU任务!

if (cpu_callback != NULL) { (*cpu_callback)(work_data); }

cudaFree(gpu parameter); //建议这个前面加入显式的同步

}

然后您将:
GPUFoo(参数);
CPUFoo(参数);
改写为:
GPUFoo(老参数,CPUFoo, CPUFoo的参数); //请根据实际原型进行修改
这么一句,即可。

以及,如果不需要有CPU的见缝插针的任务,您依然可以: GPUFoo(老参数, NULL, NULL);

您觉得呢?

感谢版主的回复。我仔细的学习过NVIDIA发布的编程手册,但是有时候遇到实际问题还是感觉模棱两可。自己的理解能得到版主的肯定,还是感觉很happy的。
另外,我这面主要是做科学计算,用cuda fortran(PGI 无限换网卡:L)。我发现貌似cuda fortran中并没有引入纹理API。是不是纹理主要还是应用于图像处理,而非科学计算?

(另外还有一个水水问题:我换了头像,但是为啥不能显示呢……)

(1)您客气了,服务您是我们的荣幸!

(2)至于fortran, 这个真心不懂的。

(3)论坛的确有换头像不变的问题,但不知道具体问题所在。不过有解决方案:按CTRL-F5强制刷新一下,即可看到您的新头像。

感谢您莅临CUDAZone China, 祝您夜安。