c++ - CUDA 动态并行，性能差-6ren

c++ - CUDA 动态并行，性能差

转载作者：塔克拉玛干更新时间：2023-11-02 23:33:30

我们在使用 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 的父内核，还调用两个子内核。第一个 child 在 [N/3, 2*N/3) 范围内 +1，第二个 child 在 [2*N/3,N) 范围内。 Childs 使用不同的流启动，因此它们可以并发。
来自主机的三个流:这个只是从 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

从图片中可以看出，主要问题是在动态并行方法中，父内核在两个子内核完成后花费过多时间关闭，这导致它需要 3X 或4 倍以上。即使考虑到最坏的情况，如果所有三个内核(父内核和两个子内核)都串行运行，也应该少得多。也就是说，每个内核有 N/3 的工作，所以整个父内核应该占用大约 3 个子内核的长度，这要少得多。 有办法解决这个问题吗？

编辑:Robert Crovella 在评论中解释了子内核以及方法 2 的序列化现象(非常感谢)。内核确实以串行方式运行的事实并不能使粗体文本中描述的问题无效(至少现在不是)。

最佳答案

调用设备运行时是“昂贵的”，就像调用主机运行时是昂贵的一样。在这种情况下，您似乎正在调用设备运行时来为每个线程创建流，即使此代码仅需要线程 0 的流。

通过修改您的代码以仅请求线程 0 的流创建，我们可以在我们为子内核启动使用单独的流的情况和我们没有为子内核使用单独的流的情况之间产生时间奇偶校验内核启动:

$ cat t370.cu
#include <stdio.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){
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if(tid == 0){
        dim3 block(BLOCKSIZE, 1, 1);
        dim3 grid( (n + BLOCKSIZE - 1)/BLOCKSIZE, 1, 1);
#ifdef USE_STREAMS
        cudaStream_t s1, s2;
        cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
        cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
        kernel_simple<<< grid, block, 0, s1 >>> (a, n, N, n);
        kernel_simple<<< grid, block, 0, s2 >>> (a, n, N, 2*n);
#else
        kernel_simple<<< grid, block >>> (a, n, N, n);
        kernel_simple<<< grid, block >>> (a, n, N, 2*n);
#endif
// these next 2 lines add noticeably to the overall timing
        cudaError_t err = cudaGetLastError();
        if (err != cudaSuccess) printf("oops1: %d\n", (int)err);
    }

    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_52 -rdc=true -lcudadevrt -o t370 t370.cu
$ ./t370 16777216 1
genarray of 3*N = 50331648.......done
cudaMemcpy:Host->Device..........done
Kernel...........................done
cudaMemcpy:Device->Host..........done
checking result..................done
rtime: 6.925632 ms
$ nvcc -arch=sm_52 -rdc=true -lcudadevrt -o t370 t370.cu -DUSE_STREAMS
$ ./t370 16777216 1
genarray of 3*N = 50331648.......done
cudaMemcpy:Host->Device..........done
Kernel...........................done
cudaMemcpy:Device->Host..........done
checking result..................done
rtime: 6.673568 ms
$

虽然没有包含在上面的测试输出中，但根据我的测试，这也使 CUDA 动态并行 (CDP) 案例 (1) 与非 CDP 案例 ( 0，2)。请注意，我们可以通过放弃调用父内核(我添加到您的代码中)中的 cudaGetLastError() 来将上述时间缩短大约 1 毫秒 (!)。

关于c++ - CUDA 动态并行，性能差，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/45201062/

文章推荐： c - 使用 getline 当行为空时退出循环

文章推荐： android - 在基于服务的类中设置监听器

文章推荐： c - 更新 select() 调用中的超时值

文章推荐： c++ - gsl::multi_span 的用途是什么？

R同时运行2个不同的代码块(并行)
有没有办法同时运行 2 个不同的代码块。我一直在研究 R 中的并行包，它们似乎都基于在循环中运行相同的函数。我正在寻找一种同时运行不同函数的方法(循环的 1 次迭代)。例如，我想在某个数据对象上创建一
C# 并行.For
无论如何增加 Parallel.For 启动后的循环次数？示例如下: var start = 0; var end = 5; Parallel.For(start, end, i => { C
Golang 并行
我是 Golang 的新手，正在尝试了解并发和并行。我阅读了下面提到的关于并发和并行的文章。我执行了相同的程序。但没有得到相同的(混合字母和字符)输出。首先获取所有字母，然后获取字符。似乎并发不工作，
R 同时循环遍历两个或多个向量 - 并行
我正在寻找同时迭代 R 中两个或多个字符向量/列表的方法，例如。有没有办法做这样的事情: foo <- c('a','c','d') bar <- c('aa','cc','dd') for(i in
Raku 并行/函数方法
我对 Raku 很陌生，我对函数式方法有疑问，尤其是 reduce。我最初有这样的方法: sub standardab{ my $mittel = mittel(@_); my $foo =
multithreading - 并行/多处理器音频处理资源
我最近花了很多时间来学习实时音频处理的细节，我发现的大多数库/工具都是c / c++代码或脚本/图形语言的形式，并在其中编译了c / c++代码。引擎盖。使用基于回调的API，与GUI或App中的其
Jmeter for循环http请求-并行
我正在使用 JMeter 进行图像负载测试。我有一个图像名称数组并遍历该数组，我通过 HTTP 请求获取所有图像。 -> loop_over_image - for loop controller
python 并行 Popen
我整个晚上都在困惑这个问题...... makeflags = ['--prefix=/usr','--libdir=/usr/lib'] rootdir='/tmp/project' ps = se
c# - 并行 C#
我正在尝试提高计算图像平均值的方法的性能。为此，我使用了两个 For 语句来迭代所有图像，因此我尝试使用一个 Parallel For 来改进它，但结果并不相同。我做错了吗？或者是什么导致了差异？
c++ - 并行 for 循环体的最佳大小
假设您有一个并行 for 循环实现，例如ConcRT parallel_for，将所有工作放在一个 for 循环体内总是最好的吗？举个例子: for(size_t i = 0; i < size()
C#并行，如何设置线程数
我想并行运行一部分代码。目前我正在使用 Parallel.For 如何让10、20或40个线程同时运行我当前的代码是: Parallel.For(1, total, (ii) =>
Paypal自适应支付(并行)550001错误
我使用 PAY API 进行了 PayPal 自适应并行支付，其中无论用户(买家)购买什么，都假设用户购买了总计 100 美元的商品。在我的自适应并行支付中，有 2 个接收方:Receiver1 和
algorithm - 并行/集群中图形节点分组的有效算法
我正在考虑让玩家加入游戏的高效算法。由于会有大量玩家，因此算法应该是异步的(即可扩展到集群中任意数量的机器)。有细节:想象有一个无向图(每个节点都是一个玩家)。玩家之间的每条边意味着玩家可以参加同一场
Java 并行 volatile i++
我有一个全局变量 volatile i = 0; 和两个线程。每个都执行以下操作: i++; System.out.print(i); 我收到以下组合。 12、21 和 22。我理解为什么我没有得到
c++ - OpenMp 并行
我有以下称为 pgain 的方法，它调用我试图并行化的方法 dist: /***************************************************************
Ruby 并行/多线程编程来读取巨大的数据库
我有一个 ruby 脚本读取一个巨大的表(约 2000 万行)，进行一些处理并将其提供给 Solr 用于索引目的。这一直是我们流程中的一大瓶颈。我打算在这里加快速度，我想实现某种并行性。我对 Ru
Golang 和并发/并行
我正在研究 Golang 并遇到一个问题，我已经研究了几天，我似乎无法理解 go routines 的概念以及它们的使用方式。基本上我是在尝试生成数百万条随机记录。我有生成随机数据的函数，并将创建一
for-loop - 并行 For 循环
我希望 for 循环使用 go 例程并行。我尝试使用 channel ，但没有用。我的主要问题是，我想在继续之前等待所有迭代完成。这就是为什么在它不起作用之前简单地编写 go 的原因。我尝试使用 ch
haskell - 并行 IO 导致终端出现随机文本输出
我正在使用 import Control.Concurrent.ParallelIO.Global main = parallel_ (map processI [1..(sdNumber runPa
R:并行 makePSOCKcluster 挂起
我正在尝试通过 makePSOCKcluster 连接到另一台计算机: library(parallel) cl ... doTryCatch -> recvData -> makeSOCKm

塔克拉玛干

个人简介

我是一名优秀的程序员,十分优秀！

作者热门文章

滴滴打车优惠券免费领取

全站热门文章

首页

博学

6Ren·AI

商城

c++ - CUDA 动态并行，性能差