【发布时间】：2017-12-25 08:17:04
【问题描述】：
我们在使用 CUDA 动态并行时遇到了性能问题。目前，CDP 的执行速度至少比传统方法慢 3 倍。 我们制作了最简单的可重现代码来显示此问题，即将数组的所有元素的值增加 +1。即，
a[0,0,0,0,0,0,0,.....,0] --> kernel +1 --> a[1,1,1,1,1,1,1,1,1]
这个简单示例的目的只是看看 CDP 是否可以像其他人一样执行，或者是否存在严重的开销。
代码在这里：
#include <stdio.h>
#include <cuda.h>
#define BLOCKSIZE 512
__global__ void kernel_parent(int *a, int n, int N);
__global__ void kernel_simple(int *a, int n, int N, int offset);
// N is the total array size
// n is the worksize for a kernel (one third of N)
__global__ void kernel_parent(int *a, int n, int N){
cudaStream_t s1, s2;
cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid == 0){
dim3 block(BLOCKSIZE, 1, 1);
dim3 grid( (n + BLOCKSIZE - 1)/BLOCKSIZE, 1, 1);
kernel_simple<<< grid, block, 0, s1 >>> (a, n, N, n);
kernel_simple<<< grid, block, 0, s2 >>> (a, n, N, 2*n);
}
a[tid] += 1;
}
__global__ void kernel_simple(int *a, int n, int N, int offset){
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int pos = tid + offset;
if(pos < N){
a[pos] += 1;
}
}
int main(int argc, char **argv){
if(argc != 3){
fprintf(stderr, "run as ./prog n method\nn multiple of 32 eg: 1024, 1048576 (1024^2), 4194304 (2048^2), 16777216 (4096^2)\nmethod:\n0 (traditional) \n1 (dynamic parallelism)\n2 (three kernels using unique streams)\n");
exit(EXIT_FAILURE);
}
int N = atoi(argv[1])*3;
int method = atoi(argv[2]);
// init array as 0
int *ah, *ad;
printf("genarray of 3*N = %i.......", N); fflush(stdout);
ah = (int*)malloc(sizeof(int)*N);
for(int i=0; i<N; ++i){
ah[i] = 0;
}
printf("done\n"); fflush(stdout);
// malloc and copy array to gpu
printf("cudaMemcpy:Host->Device..........", N); fflush(stdout);
cudaMalloc(&ad, sizeof(int)*N);
cudaMemcpy(ad, ah, sizeof(int)*N, cudaMemcpyHostToDevice);
printf("done\n"); fflush(stdout);
// kernel launch (timed)
cudaStream_t s1, s2, s3;
cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
cudaStreamCreateWithFlags(&s3, cudaStreamNonBlocking);
cudaEvent_t start, stop;
float rtime = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
printf("Kernel...........................", N); fflush(stdout);
if(method == 0){
// CLASSIC KERNEL LAUNCH
dim3 block(BLOCKSIZE, 1, 1);
dim3 grid( (N + BLOCKSIZE - 1)/BLOCKSIZE, 1, 1);
cudaEventRecord(start, 0);
kernel_simple<<< grid, block >>> (ad, N, N, 0);
cudaDeviceSynchronize();
cudaEventRecord(stop, 0);
}
else if(method == 1){
// DYNAMIC PARALLELISM
dim3 block(BLOCKSIZE, 1, 1);
dim3 grid( (N/3 + BLOCKSIZE - 1)/BLOCKSIZE, 1, 1);
cudaEventRecord(start, 0);
kernel_parent<<< grid, block, 0, s1 >>> (ad, N/3, N);
cudaDeviceSynchronize();
cudaEventRecord(stop, 0);
}
else{
// THREE CONCURRENT KERNEL LAUNCHES USING STREAMS
dim3 block(BLOCKSIZE, 1, 1);
dim3 grid( (N/3 + BLOCKSIZE - 1)/BLOCKSIZE, 1, 1);
cudaEventRecord(start, 0);
kernel_simple<<< grid, block, 0, s1 >>> (ad, N/3, N, 0);
kernel_simple<<< grid, block, 0, s2 >>> (ad, N/3, N, N/3);
kernel_simple<<< grid, block, 0, s3 >>> (ad, N/3, N, 2*(N/3));
cudaDeviceSynchronize();
cudaEventRecord(stop, 0);
}
printf("done\n"); fflush(stdout);
printf("cudaMemcpy:Device->Host..........", N); fflush(stdout);
cudaMemcpy(ah, ad, sizeof(int)*N, cudaMemcpyDeviceToHost);
printf("done\n"); fflush(stdout);
printf("checking result.................."); fflush(stdout);
for(int i=0; i<N; ++i){
if(ah[i] != 1){
fprintf(stderr, "bad element: a[%i] = %i\n", i, ah[i]);
exit(EXIT_FAILURE);
}
}
printf("done\n"); fflush(stdout);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&rtime, start, stop);
printf("rtime: %f ms\n", rtime); fflush(stdout);
return EXIT_SUCCESS;
}
可以编译
nvcc -arch=sm_35 -rdc=true -lineinfo -lcudadevrt -use_fast_math main.cu -o prog
这个例子可以用 3 种方法计算结果：
- 简单内核：只需对阵列进行一次经典内核 +1 传递。
- 动态并行：从 main() 调用在 [0,N/3) 范围内 +1 的父内核，同时调用两个子内核。第一个孩子在 [N/3, 2*N/3) 范围内执行 +1，第二个孩子在 [2*N/3,N) 范围内执行 +1。子进程使用不同的流启动，因此它们可以并发。
- 来自主机的三个流：这个仅从 main() 启动三个非阻塞流，一个用于数组的三分之一。
我得到了方法 0（简单内核）的以下配置文件： 方法1（动态并行）的以下内容： 以下是方法 2（来自主机的三个流） 运行时间是这样的：
➜ simple-cdp git:(master) ✗ ./prog 16777216 0
genarray of 3*N = 50331648.......done
cudaMemcpy:Host->Device..........done
Kernel...........................done
cudaMemcpy:Device->Host..........done
checking result..................done
rtime: 1.140928 ms
➜ simple-cdp git:(master) ✗ ./prog 16777216 1
genarray of 3*N = 50331648.......done
cudaMemcpy:Host->Device..........done
Kernel...........................done
cudaMemcpy:Device->Host..........done
checking result..................done
rtime: 5.790048 ms
➜ simple-cdp git:(master) ✗ ./prog 16777216 2
genarray of 3*N = 50331648.......done
cudaMemcpy:Host->Device..........done
Kernel...........................done
cudaMemcpy:Device->Host..........done
checking result..................done
rtime: 1.011936 ms
从图片中可以看出，主要问题是在动态并行方法中，父内核在两个子内核完成后花费了过多的时间来关闭，这导致它需要 3 倍或4 倍以上。 即使考虑最坏的情况，如果所有三个内核（父内核和两个子内核）都串行运行，它应该花费更少。即，每个内核有 N/3 的工作，因此整个父内核应该花费大约 3 个子内核，这要少得多。 有没有办法解决这个问题？
编辑：Robert Crovella 在 cmets 中解释了子内核的序列化现象以及方法 2（非常感谢）。内核确实以串行方式运行这一事实不会使以粗体文本描述的问题无效（至少现在不是）。
【问题讨论】：
-
关于序列化，序列化是由于内核的大小。完全占用 GPU 的内核启动将完全占用 GPU，并防止后续内核占用 GPU。实际上，在实践中很难目睹并发内核执行。研究相关的 CUDA 示例代码，您会发现该内核经过精心设计，可以限制 GPU 资源的使用以实现并发。如果您想了解内核并发性，请运行 CUDA 示例代码，并学习如何设计类似的代码。
-
我明白了。对于这个例子，设计几乎没有任何意义，但我理解你的解释，因此我不会认为这是一个问题，因为 GPU 正在满负荷或接近满负荷工作。在我的真实示例中，递归继续作为二叉树进行，因此最终我确实生成了将从并发中受益的小内核。然后我将不得不关注第二个问题，这是导致 3 倍或更多减速的原因。会不会是同一个原因，也就是父亲占据了GPU？但是这三个方面的工作都是 N/3，所以没有理由花更多时间在工作上
-
我并不是说我理解关于动态并行性将父内核的持续时间延长那么长的报告。我暂时无法解释。连载对我来说并不奇怪，但它似乎（从我的角度来看）是两个问题中较小的一个。换句话说，我同意。但我还没有调查其他问题。第一步是尝试重现它并稍微研究一下您的代码。
-
谢谢，如果您有机会重现代码并发回您的发现，这将非常有用。
-
似乎“cudaStreamCreateWithFlags(...)”函数是产生额外时间的函数。在完全没有流的情况下启动（也没有创建）使得 CDP 至少在非分析时间内几乎与其他运行一样快。这样做的坏处是我们消除了并发内核的任何机会，这对于最终会在某些时候产生小内核的递归算法来说是一个糟糕的打击。现在，如果我们从不同的线程块启动内核，如果一个内核的利用率很低，那会允许并发的机会吗？
标签： c++ cuda dynamic-parallelism cuda-streams