CUDA编程--内核嵌套

写在前面的话：很多算法中都需要内核嵌套，刚学CUDA可能会比较乱，所以这里我进行一个记录。（内核嵌套无法兼容Cmakelist，目前也没有特别好的办法，只能用命令行进行编译）--更新：已解决内核嵌套无法使用Cmakelist问题，将在后面文章进行阐述

本内容想阐述清楚以下内容：

1. 如何进行内核嵌套，代码结构是啥样子
2. 内核嵌套时线程号如何计算

话不多说，我们直接用代码来进行说明，先阐述两种内核嵌套的形式：内核内调用另外一个内核， 内核内调用自己。

内核内调用另外一个内核：

创建KernelNesting.cu文件，并编写以下代码：

#include <stdio.h>
__global__ void  sub_kernel( )
{
    int th_index = blockIdx.x*blockDim.x + threadIdx.x;
    printf("-------> sub_kernel thread number: %d \n", th_index);
}

__global__ void  kernel( )
{
    int th_index = blockIdx.x*blockDim.x + threadIdx.x;
    printf("-------> kernel thread number: %d \n", th_index);
    sub_kernel<<<2,2>>>();
}

int main(void )
{
    kernel<<<2,2>>>();
    cudaDeviceReset();
    return 0;
}

使用以下命令进行编译：

 nvcc -arch sm_75 -rdc=true KernelNesting.cu -o kernelnesting

这里我们新增了-rdc=true， 这表示支持循环嵌套， 其他参数的意义可以参考这篇文章：实战:Hello World——CUDA_MacalDan的博客-CSDN博客

编译成功后，我们可以看到一个名为kernelnesting的可执行文件，通过以下命令运行

./kernelnesting

结果如下：

不难发现，内核嵌套后，新的内核的线程依旧从0开始计算，即我们每开启一个内核其线程都是从0开始。 

内核内调用自己：

我们在上面的文件中稍作修改，仅修改一个函数即可，如下：

__global__ void  kernel( )
{
    int th_index = blockIdx.x*blockDim.x + threadIdx.x;
    printf("-------> kernel thread number: %d \n", th_index);
//    sub_kernel<<<2,2>>>();
    kernel<<<2,2>>>();
}

修改后我们也能预测出结果来了，该文件会一直调用自己，永远不会停止。。。

同样我们编译并运行文件

nvcc -arch sm_75 -rdc=true KernelNesting.cu -o kernelnesting_ziji

./kernelnesting_ziji

得出结果如下：

程序会一致跑下去，永远不停止，除非你点个赞，然后我改下程序（加个停止条件）。其次，我们可以看出，自己调用自己时线程号也会归零， 即我们每开启一个内核其线程都是从0开始。 

最后，完整代码以及其他demo可以在我的github中获取（不定期更新）：

https://github.com/weiguangzhao/cuda_demohttps://github.com/weiguangzhao/cuda_demo