原标题：CUDA 动态并行

原文来自：CSDN      原文链接：https://blog.csdn.net/weixin_36709051/article/details/103807551

CUDA 动态并行

1.cuda执行模型

    最近开始学习cuda，在研读完GPU硬件结构与CUDA编程接口后，对cuda的理解如下：cuda执行过程分为五个部分。1.分配主机内存与设备显存；2. 将数据从内存复制到显存；3.执行核函数；4.将数据从显存复制到内存；5.释放主机内存与设备显存。

    核函数kernel<<<grid,block>>>是cuda上执行的核心程序，它有两个层次网格(grid)与线程块(block)，由一个核函数启动的所有线程称为一个线程块，线程块内又有很多线程。通过线程块索引与块内的线程索引我们可以很方便的管理线程。gird与block都是dim3类型，总共有三个维度，每个维度类型为size_t，默认为1。

2.实际编程遇到的问题

    cuda编程从概念上看起来非常容易，但在实际操作过程中却遇到了很多的麻烦。通常每个线程的数据是从全局内存(global memory)中获取的，但对于一些复杂的操作，每个线程需要开辟自己独立的空间来对数据进行处理。比如在CFAR恒虚警率检测器算法中，每一个线程要对滑动窗内获得的数据进行统计分布拟合，显然要将全局内存中的图像数据复制到线程自身的堆栈空间中来进一步处理。

    这个问题困扰了我很久，cuda不支持动态内存分配，不像c++那样，可以在程序运行过程中动态的在堆上new一段内存空间。难道要提前分配好每一个线程的显存空间，然后根据块索引与线程索引来找到自身对应的显存空间？感觉这样做十分的笨拙与不靠谱，我开始在度娘与google上搜索CUDA动态分配显存，发现基本没有相关信息，有的说可以在核函数内new与delete，我尝试了一下，确实可以，但是每个线程分配的double数目超过500个后，就会报错，另外这段空间应该是在内存上分配的而不是在显存上分配的。所以这个方案很快被否定了。接下来就讲一下一个可行的解决方案，CUDA动态并行。

3.CUDA动态并行

    了解cuda动态并行是从做排序开始的，在gpu上需要对一段数组进行排序，开始写了一份选择排序，但显然O(n²)的复杂度是无法接受的。在CUDA的实例samples/0_simple/cdpSimpleQuicksort/中提供了一段快排程序， 部分核心代码如下：

__global__ void cdp_simple_quicksort(unsigned int *data, int left, int right, int depth)
{
    ...
    ...
    if (left < (rptr-data))
    {
        cudaStream_t s;
        cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking);
        cdp_simple_quicksort<<< 1, 1, 0, s >>>(data, left, nright, depth+1);
        cudaStreamDestroy(s);
    }
    // Launch a new block to sort the right part.
    if ((lptr-data) < right)
    {
        cudaStream_t s1;
        cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking);
        cdp_simple_quicksort<<< 1, 1, 0, s1 >>>(data, nleft, right, depth+1);
        cudaStreamDestroy(s1);
    }
}

3.1 选择设备计算能力

    在这段程序中竟然有核函数的递归调用，在我之前的认知中，核函数与设备函数是不支持递归调用的，于是我尝试自己编译这段代码，而非该目录下的makefile，报错如下。

nvcc -I/usr/local/cuda/samples/common/inc -c cdpSimpleQuicksort.cu
(其中-I为包含helper_cuda.h的路径)

cdpSimpleQuicksort.cu(114): error: calling a __global__ function("cdp_simple_quicksort") from a __global__ function("cdp_simple_quicksort") is only allowed on the compute_35 architecture or above

    在该报错信息中我们发现核函数调用核函数只能在compute_35或更高的架构上，经过一番查找发现nvcc默认的编译架构是sm_30，其只支持一些基本的特性。不同计算架构具有的功能如下表，来自nvidia官网。我使用的GPU是GTX 1080 Ti， 其计算能力6.1，满足3.5要求，所以问题肯定出在了编译上。

3.2 选择独立编译模式

    在nvidia官网https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html中详细的介绍了nvcc的使用方法，–gpu-architecture（-arch）来确定文件编译的GPU架构。在添加了编译架构后，之前的核函数调用报错信息消失了，但还是有报错如下：

nvcc -arch=sm_50  -I/usr/local/cuda/samples/common/inc -c cdpSimpleQuicksort.cu

error: kernel launch from __device__ or __global__ functions requires separate compilation mode

    报错信息：从设备函数中或核函数中启动核函数需要独立编译模式。在nvidia官网的介绍中，–relocatable-device-code(-rdc)，允许产生可重定位的设备代码，如果设置为false，则生成可执行设备代码。若设置为true生成可重定位设备代码必须被连接才可执行。

3.3 可重定位设备目标文件生成主机可连接的目标文件

加入-rdc=true后生成的是可重定位设备目标文件，必须要生成主机可连接的可执行设备目标文件。nvcc选项–device-link(-dlink)。在生成设备目标文件后，由g++生成可执行文件。

nvcc -arch=sm_50 -rdc=true  -I/usr/local/cuda/samples/common/inc -c cdpSimpleQuicksort.cu
(生成可重定位设备目标文件)
nvcc -arch=sm_50 -dlink cdpSimpleQuicksort.o -o device_link.o
(生成主机可连接目标文件)
g++ device_link.o cdpSimpleQuicksort.o -o Qsort.out -L/usr/local/cuda/lib64 -lcudart -lcudadevrt
(g++生成主机可执行文件)

若没有-dlink的过程，在生成可执行文件过程中将会出现如下的报错信息。

 undefined reference to `__cudaRegisterLinkedBinary_53_tmpxft_00003838_00000000_6_cdpSimpleQuicksort_cpp1_ii_e4d37a34'
collect2: error: ld returned 1 exit status

4. CUDA快速排序算法

    近期在做CFAR舰船检测相关CUDA加速算法，对于CA-CFAR单元平均虚警率算法，需要对杂波像素做升序排序，并取前70%的像素计算均值，相关github代码连接:https://github.com/sty16/ship-detection，选择动态并行后可以支持设备函数的递归调用，一个简单的快速排序算法如下：

__device__  void simple_quicksort(double *data, int left, int right, int depth)
{
    if (depth >= MAX_DEPTH || right-left <= INSERTION_SORT)
    {
        selection_sort(data, left, right);
        return;
    }
    double *lptr = &data[left];
    double *rptr = &data[right];
    double  pivot = data[(left+right)/2];
    while(lptr <= rptr)
    {
        double lval = *lptr;
        double rval = *rptr;
        while(lval < pivot)
        {
            lptr++;
            lval = *lptr;
        }
        while(rval > pivot)
        {
            rptr--;
            rval = *rptr;
        }
        if(lptr <= rptr)
        {
            *lptr++ = rval;
            *rptr-- = lval;
        }
    }
    int nright = rptr - data;
    int nleft  = lptr - data;
    if (left < (rptr-data))
    {
        simple_quicksort(data, left, nright, depth+1);
    }
    if ((lptr-data) < right)
    {
        simple_quicksort(data, nleft, right, depth+1);
    }
}

免责声明：本文来自互联网新闻客户端自媒体，不代表本网的观点和立场。

合作及投稿邮箱：E-mail:editor@tusaishared.com