如何实现nvidia显卡的cuda的多kernel并发执行？？？
相关：
CPU端多进程/多线程调用CUDA是否可以加速？？？
参考：
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）
const int N = 1 << 20; __global__ void kernel(float *x, int n) { int tid = threadIdx.x + blockIdx.x * blockDim.x; for (int i = tid; i < n; i += blockDim.x * gridDim.x) { x[i] = sqrt(pow(3.14159,i)); } } int main() { const int num_streams = 8; cudaStream_t streams[num_streams]; float *data[num_streams]; for (int i = 0; i < num_streams; i++) { cudaStreamCreate(&streams[i]); cudaMalloc(&data[i], N * sizeof(float)); // launch one worker kernel per stream kernel<<<1, 64, 0, streams[i]>>>(data[i], N); // launch a dummy kernel on the default stream kernel<<<1, 1>>>(0, 0); } cudaDeviceReset(); return 0; }
编译：
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);
修改后代码：
const int N = 1 << 20; __global__ void kernel(float *x, int n) { int tid = threadIdx.x + blockIdx.x * blockDim.x; for (int i = tid; i < n; i += blockDim.x * gridDim.x) { x[i] = sqrt(pow(3.14159,i)); } } int main() { const int num_streams = 8; cudaStream_t streams[num_streams]; float *data[num_streams]; for (int i = 0; i < num_streams; i++) { cudaStreamCreate(&streams[i]); cudaMalloc(&data[i], N * sizeof(float)); // launch one worker kernel per stream kernel<<<1, 64, 0, streams[i]>>>(data[i], N); // launch a dummy kernel on the default stream // kernel<<<1, 1>>>(0, 0); } cudaDeviceReset(); return 0; }
编译：
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队列阻塞，代码如下（与第一个代码相同）：
const int N = 1 << 20; __global__ void kernel(float *x, int n) { int tid = threadIdx.x + blockIdx.x * blockDim.x; for (int i = tid; i < n; i += blockDim.x * gridDim.x) { x[i] = sqrt(pow(3.14159,i)); } } int main() { const int num_streams = 8; cudaStream_t streams[num_streams]; float *data[num_streams]; for (int i = 0; i < num_streams; i++) { cudaStreamCreate(&streams[i]); cudaMalloc(&data[i], N * sizeof(float)); // launch one worker kernel per stream kernel<<<1, 64, 0, streams[i]>>>(data[i], N); // launch a dummy kernel on the default stream kernel<<<1, 1>>>(0, 0); } cudaDeviceReset(); return 0; }
编译命令：
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队列来运行操作，那么效果如何呢？
给出代码：
#include <pthread.h> #include <stdio.h> const int N = 1 << 20; __global__ void kernel(float *x, int n) { int tid = threadIdx.x + blockIdx.x * blockDim.x; for (int i = tid; i < n; i += blockDim.x * gridDim.x) { x[i] = sqrt(pow(3.14159,i)); } } void *launch_kernel(void *dummy) { float *data; cudaMalloc(&data, N * sizeof(float)); kernel<<<1, 64>>>(data, N); cudaStreamSynchronize(0); return NULL; } int main() { const int num_threads = 8; pthread_t threads[num_threads]; for (int i = 0; i < num_threads; i++) { if (pthread_create(&threads[i], NULL, launch_kernel, 0)) { fprintf(stderr, "Error creating threadn"); return 1; } } for (int i = 0; i < num_threads; i++) { if(pthread_join(threads[i], NULL)) { fprintf(stderr, "Error joining threadn"); return 2; } } cudaDeviceReset(); return 0; }
默认编译：
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是什么含义呢？
从上面的英文内容我们可以知道默认情况下每个CUDA代码在GPU上运行都会在context下有一个default stream的kernel队列，而这个default stream队列中的kernel执行会阻塞其他stream队列中的kernel操作，从而导致多个stream队列中的kernel操作无法并行。在编译的时候加入参数--default-stream per-thread，就可以使CPU端的每个线程默认调用的default stream队列映射到一个non-default stream队列中，这样就避免了因为default stream队列引起的同步阻塞。
====================================================
一个关于CUDA多流并发（多kernel并发）的PPT：
https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf
posted on 2022-11-07 12:31 Angry_Panda 阅读(2506) 评论(0) 编辑 收藏 举报
