c++ - Cuda 动态并行和递归性能

Question

我真的很好奇动态并行的性能和给你的递归能力。所以我对 cpu 递归和设备（gpu）递归做了一个简单的基准测试：

#include <stdio.h>
#include <curand.h>
#include <cuda_runtime_api.h>
#include <time.h>
#include <sys/time.h>

__global__ void recurse(int level)
{
    if (level<0)
    {
        return;
    }
    level--;
    recurse<<<1,1>>>(level);
    recurse<<<1,1>>>(level);
    return;
}

void serial_recurse(int level)
{
    if (level<0)
    {
        return;
    }
    level--;
    serial_recurse(level);
    serial_recurse(level);
    return;

}


int main()
{
    struct timeval startwtime,endwtime;
    double seq_time;

    printf("Startin here\n");

    int level_1 = 22;
    gettimeofday(&startwtime,NULL);

    serial_recurse(level_1);

    gettimeofday(&endwtime,NULL);

    seq_time = (double)((endwtime.tv_usec - startwtime.tv_usec)/1.0e6+endwtime.tv_sec - startwtime.tv_sec);
    printf("serial time = %f \n", seq_time);

    int level=22;
    
    gettimeofday(&startwtime,NULL);

    recurse<<<1,1>>>(level);
    cudaDeviceSynchronize();

    gettimeofday(&endwtime,NULL);

    seq_time = (double)((endwtime.tv_usec - startwtime.tv_usec)/1.0e6+endwtime.tv_sec - startwtime.tv_sec);
    printf("parralel time = %f \n", seq_time);

    return 0;
}

结果对于动态并行性来说确实令人失望：
cpu: 0.04666500 sec
gpu: 283.620204 sec

有什么方法可以通过动态并行获得更好的性能，或者还为时过早？

标签： c++ccuda