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

coder 2024-02-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 的父内核，还调用两个子内核。第一个 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/

amp 43 int BLOCKSIZE kernel c++cuda dynamic-parallelism cuda-streams

有关c++ - CUDA 动态并行，性能差的更多相关文章

ruby-on-rails - 如何优雅地重启 thin + nginx？ - 2
我的瘦服务器配置了nginx，我的ROR应用程序正在它们上运行。在我发布代码更新时运行thinrestart会给我的应用程序带来一些停机时间。我试图弄清楚如何优雅地重启正在运行的Thin实例，但找不到好的解决方案。有没有人能做到这一点？最佳答案 #Restartjustthethinserverdescribedbythatconfigsudothin-C/etc/thin/mysite.ymlrestartNginx将继续运行并代理请求。如果您将Nginx设置为使用多个上游服务器，例如server{listen80;server
ruby - 使用 `+=` 和 `send` 方法 - 2
如何将send与+=一起使用？a=20;a.send"+=",10undefinedmethod`+='for20:Fixnuma=20;a+=10=>30 最佳答案恐怕你不能。+=不是方法，而是语法糖。参见http://www.ruby-doc.org/docs/ProgrammingRuby/html/tut_expressions.html它说Incommonwithmanyotherlanguages,Rubyhasasyntacticshortcut:a=a+2maybewrittenasa+=2.你能做的最好的事情是:
ruby - 如何计算 Liquid 中的变量 +1 - 2
我对如何计算通过{%assignvar=0%}赋值的变量加一完全感到困惑。这应该是最简单的任务。到目前为止，这是我尝试过的:{%assignamount=0%}{%forvariantinproduct.variants%}{%assignamount=amount+1%}{%endfor%}Amount:{{amount}}结果总是0。也许我忽略了一些明显的东西。也许有更好的方法。我想要存档的只是获取运行的迭代次数。最佳答案因为{{incrementamount}}将输出您的变量值并且不会影响{%assign%}定义的变量，我
ruby - 带括号和 splat 运算符的并行赋值 - 2
我明白了:x,(y,z)=1,*[2,3]x#=>1y#=>2z#=>nil我想知道为什么z的值为nil。最佳答案 x,(y,z)=1,*[2,3]右侧的splat*是内联扩展的，所以它等同于:x,(y,z)=1,2,3左边带括号的列表被视为嵌套赋值，所以它等价于:x=1y,z=23被丢弃，而z被分配给nil。关于ruby-带括号和splat运算符的并行赋值，我们在StackOverflow上找到一个类似的问题： https://stackoverflow
ruby - 使对象的行为类似于 ruby 中并行分配的数组 - 2
假设您在Ruby中执行此操作:ar=[1,2]x,y=ar然后，x==1和y==2。是否有一种方法可以在我自己的类中定义，从而产生相同的效果？例如rb=AllYourCode.newx,y=rb到目前为止，对于这样的赋值，我所能做的就是使x==rb和y=nil。Python有这样一个特性:>>>classFoo:...def__iter__(self):...returniter([1,2])...>>>x,y=Foo()>>>x1>>>y2 最佳答案是的。定义#to_ary。这将使您的对象被视为要分配的数组。irb>o=Obje
ruby - 在 Ruby 中动态创建数组 - 2
有没有办法在Ruby中动态创建数组？例如，假设我想遍历用户输入的书籍数组:books=gets.chomp用户输入:"TheGreatGatsby,CrimeandPunishment,Dracula,Fahrenheit451,PrideandPrejudice,SenseandSensibility,Slaughterhouse-Five,TheAdventuresofHuckleberryFinn"我把它变成一个数组:books_array=books.split(",")现在，对于用户输入的每一本书，我想用Ruby创建一个数组。伪代码来做到这一点:x=0books_array.
ruby - 是否可以将 IRB 提示配置为动态更改？ - 2
我想在IRB中浏览文件系统并让提示更改以反射(reflect)当前工作目录，但我不知道如何在每个命令后进行提示更新。最终，我想在日常工作中更多地使用IRB，让bash溜走。我在我的.irbrc中试过这个:require'fileutils'includeFileUtilsIRB.conf[:PROMPT][:CUSTOM]={:PROMPT_N=>"\e[1m:\e[m",:PROMPT_I=>"\e[1m#{pwd}>\e[m",:PROMPT_S=>"FOO",:PROMPT_C=>"\e[1m#{pwd}>\e[m",:RETURN=>""}IRB.conf[:PROMPT_MO
arrays - Ruby 数组 += vs 推送 - 2
我有一个数组数组，想将元素附加到子数组。+=做我想做的，但我想了解为什么push不做。我期望的行为(并与+=一起工作):b=Array.new(3,[])b[0]+=["apple"]b[1]+=["orange"]b[2]+=["frog"]b=>[["苹果"],["橙子"],["Frog"]]通过推送，我将推送的元素附加到每个子数组(为什么？):a=Array.new(3,[])a[0].push("apple")a[1].push("orange")a[2].push("frog")a=>[[“苹果”、“橙子”、“Frog”]、[“苹果”、“橙子”、“Frog”]、[“苹果”、“
Ruby 的数字方法性能 - 2
我正在使用Ruby解决一些ProjectEuler问题，特别是这里我要讨论的问题25(Fibonacci数列中包含1000位数字的第一项的索引是多少？)。起初，我使用的是Ruby2.2.3，我将问题编码为:number=3a=1b=2whileb.to_s.length但后来我发现2.4.2版本有一个名为digits的方法，这正是我需要的。我转换为代码:whileb.digits.length当我比较这两种方法时，digits慢得多。时间./025/problem025.rb0.13s用户0.02s系统80%cpu0.190总计./025/problem025.rb2.19s用户0.0
ruby - Ruby 性能中的计时器 - 2
我正在寻找一个用ruby演示计时器的在线示例，并发现了下面的代码。它按预期工作，但这个简单的程序使用30Mo内存(如Windows任务管理器中所示)和太多CPU有意义吗？非常感谢deftime_blockstart_time=Time.nowThread.new{yield}Time.now-start_timeenddefrepeat_every(seconds)whiletruedotime_spent=time_block{yield}#Tohandle-vesleepinteravalsleep(seconds-time_spent)iftime_spent

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

有关c++ - CUDA 动态并行，性能差的更多相关文章

随机推荐