使用CUDA为GPU同时启动多个内核

Question

是否可以同时启动两个执行独立任务的内核.例如,如果我有这个Cuda代码

// host and device initialization
.......
.......

// launch kernel1
myMethod1 <<<.... >>> (params);

// launch kernel2
myMethod2 <<<.....>>> (params);

假设这些内核是独立的,是否有设施可以同时为每个内核分配几个网格/块.CUDA/OpenCL是否有此规定.

Answer 1

只有具有CUDA计算能力2.0及更高版本(即Fermi)的设备才能支持多个同步内核执行.请参阅CUDA 3.0编程指南的第3.2.6.3节,其中指出:

某些计算能力2.0的设备可以同时执行多个内核.应用程序可以通过调用cudaGetDeviceProperties()和检查concurrentKernels属性来查询此功能 .

设备可以同时执行的最大内核启动次数是四次.

来自一个CUDA上下文的内核不能与来自另一个CUDA上下文的内核同时执行.

使用许多纹理或大量本地内存的内核不太可能与其他内核并发执行.

Answer 2

并发内核需要SM 2.0或更高版本.

要获得并发执行,您需要手动指示两个内核之间没有依赖关系.这是因为编译器无法确定一个内核不会修改另一个内核正在使用的数据,这可能是通过读取和写入相同的缓冲区看起来很简单,但实际上很难检测,因为内部可能有指针数据结构等.

要表达独立性,您必须在不同的流中启动内核.triple-chevron语法中的第四个参数指定流,查看编程指南或SDK concurrentKernels示例.