参考:
https://developer.nvidia.com/blog/gpu-pro-tip-cuda-7-streams-simplify-concurrency/
====================================================
如何实现nvidia显卡的cuda的多kernel并发执行???
主要参考:GPU Pro Tip: CUDA 7 Streams Simplify Concurrency
====================================================
2022年11月11日更新
在nvidia显卡的CUDA计算中default stream是比较特殊的存在,任何没有指定的GPU上的操作都是在default stream中执行的,而default stream队列中操作的执行有一个特定就是会独占整个CPU进程在GPU端创建的context环境,也就是说default stream中的操作执行的话不论是否有其他stream队列中有操作都需要等待default stream中的操作结束才可以执行;其他non-default stream队列中如果有操作在执行,那么default stream中的操作将阻塞,直至独占整个context。如果default stream队列和non-default stream队列中都有操作,那么就会根据CPU端发送到GPU端执行命令的先后进行排队执行。
====================================================
编写多流并行(多kernel并行)的CUDA代码:(源自:GPU Pro Tip: CUDA 7 Streams Simplify Concurrency)
View Code
编译:
nvcc ./stream_test.cu -o stream_legacy
使用NVIDIA Visual Profiler (nvvp
)查看运行情况:
可以看到虽然在代码中将多个kernel的操作写在了不同的stream队列中,而且cuda代码运行的过程中也确实将不同的kernel操作放入到了不同的stream中执行,但是不同的stream的kernel并没有实现并行而是仍然串行。其主要原因就是不同的stream队列操作后都有一个default stream队列的操作,在默认的编译条件下default stream队列中的操作将阻塞其他stream队列中的操作,也是修改代码,剔除掉default stream队列中的操作:
// launch a dummy kernel on the default stream
kernel<<<1, 1>>>(0, 0);
修改后代码:
View Code
编译:
nvcc ./stream_test.cu -o stream_legacy
使用NVIDIA Visual Profiler (nvvp
)查看运行情况:
可以看到在有没有default stream队列的操作后所有其他stream队列中的kernel操作实现了并行。
如果在编译cuda代码的时候加入参数--default-stream per-thread,就可以将default stream队列的操作映射到其他stream队列中,这样就不会使其他stream队列被default stream队列阻塞,代码如下(与第一个代码相同):
View Code
编译命令:
nvcc --default-stream per-thread ./stream_test.cu -o stream_per-thread
使用NVIDIA Visual Profiler (nvvp
)查看运行情况:
可以看到加入编译参数--default-stream per-thread后所有的原先在default stream中的kernel操作都被映射到了stream 15队列中,并且stream 15队列中的kernel操作没有implicit隐式的与其他stream队列中的操作进行同步。
-----------------------------------------------------
如果同样的cuda操作使用CPU端多线程调用并且将每次的kernel调用都默认使用default stream队列来运行操作,那么效果如何呢?
给出代码:
View Code
默认编译:
nvcc ./pthread_test.cu -o pthreads_legacy
使用NVIDIA Visual Profiler (nvvp
)查看运行情况:
可以看到虽然在CPU端使用多线程调用kernel操作,但是所有的kernel操作都是使用的default stream队列,因此并不能实现多个kernel操作的GPU端并行。
如果在编译cuda代码的时候加入参数--default-stream per-thread,就可以将default stream队列的操作映射到其他stream队列中(代码与上个代码相同):
加参数编译:
nvcc --default-stream per-thread ./pthread_test.cu -o pthreads_per_thread
使用NVIDIA Visual Profiler (nvvp
)查看运行情况:
可以看到加参数编译后CPU端的每个线程调用的kernel都映射到了一个新的stream队列中,实现了GPU端的多kernel并行操作。
===========================================================
看到前面的内容可以知道,想要GPU上进行多kernel的并行需要把不同的kernel操作写在不同的stream队列中,并且一定要在编译的时候加参数:--default-stream per-thread,虽然在单进程单线程的情况下不使用default stream队列存在也可以的特例。
那么参数:--default-stream per-thread是什么含义呢?
--default-stream per-thread,就可以使CPU端的每个线程默认调用的default stream队列映射到一个non-default stream队列中,这样就避免了因为default stream队列引起的同步阻塞。
====================================================
一个关于CUDA多流并发(多kernel并发)的PPT:
https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf