前言

杜老师推出的 tensorRT从零起步高性能部署 课程，之前有看过一遍，但是没有做笔记，很多东西也忘了。这次重新撸一遍，顺便记记笔记。

本次课程学习精简 CUDA 教程-核函数

课程大纲可看下面的思维导图

1. 核函数

关于核函数你需要知道：

1. 核函数是 cuda 编程的关键

2. 通过 xxx.cu 创建一个 cudac 程序文件，并把 cu 交给 nvcc 编译，才能识别 cuda 语法

   • nvcc 是 nvidia 的一个 c++ 编译器，是用来编译 cudac 程序的

3. __global__ 表示为核函数，由 host 调用。

4. __deivce__ 表示为设备函数，由 device 调用

5. __host__ 表示为主机函数，由 host 调用。__shared__ 表示变量为共享变量

   • 一个函数可以既是设备函数又是主机函数，可以同时被 __device__ __host__ 修饰

6. host 调用核函数：function<<<gridDim，blockDim，sharedMemorySize，stream>>>(args…)

   • stream 是上节课提到的流，在进行异步管理的时候可以控制它。sharedMemorySize 是共享内存的大小
   • gridDim 和 blockDim 用于告诉核函数该启动多少个线程，二者都是内置变量，其变量类型是 dim3
   • 启动的总线程数量 nthreads = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z
   • gridDim 和 blockDim 都是有约束的，可以通过 runtime API 或者 deviceQuery 进行查询。gridDims(21亿，65536，65536)，blockDim(1024，64，64) blockDim.x * blockDim.y * blockDim.z <= 1024

7. 只有 __global__ 修饰的函数才可以 <<<>>> 的方式调用

8. 调用核函数是传值的，不能传引用，可以传递类、结构体等，核函数可以是模板，返回值必须是 void

9. 核函数的执行，是异步的，也就是立即返回的

10. 线程 layout 主要用到 blockDim、gridDim

11. 核函数内访问线程索引主要用到 threadIdx、blockIdx、blockDim、gridDim 这些内置变量

我们之前有提到将 host 即 CPU 上的数据拷贝到 device 即 GPU 上，目的是什么呢？目的当然是利用 GPU 的高性能并行计算能力，那具体怎么在 GPU 上利用这些数据来完成指定的计算呢？这就需要你来调用 CUDA 中的核函数 (kernel) 来执行并行计算。

kernel 是 CUDA 编程中一个重要的概念，指的是在 device 上线程并行执行的函数，核函数使用 __global__ 符号声明，在调用时使用 <<<grid, block>>> 来指定核函数 kernel 要执行的线程数量，在 CUDA 中的每一个线程都要执行核函数，并且每个线程会分配一个唯一的 线程号 thread ID，这个 ID 值可以通过核函数内置变量 threadIdx 来获得。

由于 GPU 实际上是异构模型，所以需要区分 host 和 device 上的代码，在 CUDA 中我们是通过函数类型限定词来区分 host 和 device 上的函数，主要有三个函数类型限定词：

• __global__ 表示核函数，在 device 上执行，由 host 调用，返回类型必须是 void
• __device__ 表示设备函数，在 device 上执行，仅可以从 device 调用
• __host__ 表示主机函数，仅可以从 host 上调用

要深刻理核函数，必须要对其的线程层级结构有一个清晰的认识。

首先 GPU 上有很多并行化的轻量级线程，kernel 在 device 上执行时实际上是启动了很多线程，一个 kernel 所启动的所有线程称为一个网格（grid），同一个网格上的所有线程共享相同的全局内存空间，grid 是线程结构的第一层次。而网格 grid 又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。

线程两层组织结构如 图1-1 所示，从图中可以看出这是一个 grid 和 block 均为 2-dim 的线程组织。那 2-dim 又是什么意思呢？这就不得不提 grid 和 block 变量类型了，grid 和 block 其实都是定义在 dim3 类型的变量，而 dim3 可以看成是包含三个无符号整数 (x,y,z) 成员的结构体变量，在定义时，缺省值初始化为 1。

图1-1 Kernel上的两层线程组织结构(2-dim)

因此 grid 和 block 可以灵活地定义为 1-dim，2-dim 以及 3-dim 结构，正常 2-dim 线程结构是比较常用的，对于 图1-2 的线程组织结构而言，grid 和 block 的定义可以如下：

dim3 grid(3, 2);
dim3 block(5, 3);
kernel_func<<<grid, block>>>(params, ...);

值得注意的是，核函数在调用时必须通过执行配置 <<<gird,block>>> 来指定 kernel 所使用的线程数及线程结构。

所以，一个线程需要两个内置变量（blockIdx，threadIdx）来唯一标识，它们都是 dim3 类型的变量，其中 blockIdx 指明了该线程在 grid 网格中的位置，而 threadIdx 指明了该线程在 block 中的位置，在 图xxx 中的 Thread(1,1) 满足：

threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

有时候，我们还想要知道一个线程在线程块（block）中的全局 ID，此时就必须还要知道 block 的组织结构，这是通过线程的内置变量 blockDim 来获得。它获取线程块（block）各个维度的大小，对于一个 2-dim 的 block($D_x,D_y$)，线程 ($x,y$) 的 ID 值为 ($x+y\ast D_x$)，如果是 3-dim 的 block($D_x,D_y,D_z$)，线程 ($x,y,z$)的 ID 值为 ($x+y\ast D_x+z\ast D_x\ast D_y$)。另外线程还有内置变量 gridDim，用于获得网格块各个维度的大小。

如果你还想要知道当前线程在所有线程中的即网格（grid）中的全局 ID，我们就需要同时用到 gridDim 和 blockDim，根据杜老师的方法可以很简单的计算出对应的全局 ID，具体如下图所示：

图1-2 网格中线程的全局ID计算

在核函数里，可以把 blockDim、gridDim 看作 shape，把 threadIdx、blockIdx 看作 index，对于全局索引的计算有个方便的记忆办法就是左乘右加，之后无论 tensor 维度有多复杂，这个方法都适用。而线程的全局索引通常会映射到指针的偏移量上，方便我们后续的操作

我们拿个简单的例子来说明，假设 grid(2,1,1) blockDim(1,1,10)

图1-3 全局索引计算示例

按照左乘右加准则，则 idx = blockIdx.x * blockDim.x + threadIdx.x

一个线程块（block）上的线程是放在同一个**流式多处理器（streaming Multi-processor，SM）**上的，但是单个 SM 的资源是有限的，这导致线程块（block）中的线程数是有限的，现代 GPU 的线程数可支持的线程数可达 1024 个。

kernel 核函数在执行时实际上会启动很多线程，这些线程在逻辑上是并行的，但是在物理层不一定。GPU 硬件的一个核心组件就是 SM，当一个 kernel 核函数被执行时，它的 grid 中的线程块就被分配到 SM 上。一个线程块只能在一个 SM 上被调度，而一个 SM 一般可以调度多个线程块，这要看 SM 本身能力。一个 kernel 的各个线程块被分配到多个 SM，因此 grid 只是逻辑层，而 SM 才是执行的物理层，如图1-4所示。

SM 的基本执行单元是线程束（warps），线程束包含 32 个线程，但是一个 SM 的同时并发的线程束数是有限的。总之，就是网格和线程块只是逻辑划分，一个 kernel 的所有线程其实在物理层不一定同时并发的。所以，kernel 的 grid 和 block 的配置不同，性能会出现差异。还要注意，由于 SM 的基本执行单元是包含 32 个线程的线程束，所以 block 大小一般要设置为 32 的倍数。

图1-4 CUDA编程的逻辑层和物理层

2. 核函数案例

核函数案例的 main.cpp 示例代码如下：

#include <cuda_runtime.h>
#include <stdio.h>

#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)

bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}

void test_print(const float* pdata, int ndata);

int main(){
    float* parray_host = nullptr;
    float* parray_device = nullptr;
    int narray = 10;
    int array_bytes = sizeof(float) * narray;

    parray_host = new float[narray];
    checkRuntime(cudaMalloc(&parray_device, array_bytes));

    for(int i = 0; i < narray; ++i)
        parray_host[i] = i;
    
    checkRuntime(cudaMemcpy(parray_device, parray_host, array_bytes, cudaMemcpyHostToDevice));
    test_print(parray_device, narray);
    checkRuntime(cudaDeviceSynchronize());

    checkRuntime(cudaFree(parray_device));
    delete[] parray_host;
    return 0;
}

核函数案例的 kernel.cu 示例代码如下：

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void test_print_kernel(const float* pdata, int ndata){

    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    /*    dims                 indexs
        gridDim.z            blockIdx.z
        gridDim.y            blockIdx.y
        gridDim.x            blockIdx.x
        blockDim.z           threadIdx.z
        blockDim.y           threadIdx.y
        blockDim.x           threadIdx.x

        Pseudo code:
        position = 0
        for i in 6:
            position *= dims[i]
            position += indexs[i]
    */
    printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

void test_print(const float* pdata, int ndata){

    // <<<gridDim, blockDim, bytes_of_shared_memory, stream>>>
    test_print_kernel<<<1, ndata, 0, nullptr>>>(pdata, ndata);

    // 在核函数执行结束后，通过cudaPeekAtLastError获取得到的代码，来知道是否出现错误
    // cudaPeekAtLastError和cudaGetLastError都可以获取得到错误代码
    // cudaGetLastError是获取错误代码并清除掉，也就是再一次执行cudaGetLastError获取的会是success
    // 而cudaPeekAtLastError是获取当前错误，但是再一次执行 cudaPeekAtLastError 或者 cudaGetLastError 拿到的还是那个错
    // cuda的错误会传递，如果这里出错了，不移除。那么后续的任意api的返回值都会是这个错误，都会失败
    cudaError_t code = cudaPeekAtLastError();
    if(code != cudaSuccess){    
        const char* err_name    = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("kernel error %s:%d  test_print_kernel failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, err_name, err_message);   
    }
}

运行效果如下：

图2-1 kernel案例运行效果

这个案例展示了如何在 CUDA 中使用核函数进行并行计算。

test_print_kernel 是一个 __global__ 修饰符标记的核函数，它将在 GPU 上执行，并由 host 调用。核函数的作用是打印传入数据数组的每个元素的值以及线程索引、块索引和块大小等信息。test_print 为主机函数负责调用核函数，<<<1, ndata, 0, nullptr>>> 是启动核函数的语法，其中 1 是块（block）的数量，ndata 是每个块中的线程（thread）数量，0 表示共享内存大小，nullptr 表示使用默认的流（stream）。

在核函数执行结束后，使用 cudaPeekAtLastError 检查是否有错误发生。如果有错误，将打印错误代码和消息。值得注意的是 cudaPeekAtLastError 和 cudaGetLastError 都可以获取得到错误代码，cudaGetLastError 是获取错误代码并清除掉，也就是再一次执行 cudaGetLastError 获取的会是 success。而 cudaPeekAtLastError 是获取当前错误，但是再一次执行 cudaPeekAtLastError 或者 cudaGetLastError 拿到的还是那个错误。cuda 的错误会传递，如果这里出错了，不移除，那么后续的任意 api 的返回值都会是这个错误，都会失败。

通过这个案例，可以了解如何定义和启动核函数，并使用线程索引、块索引和块大小等信息来实现并行计算。在实际应用中，可以根据需要编写更复杂的核函数来处理实际计算任务。

关于核函数的知识点如下：(from 杜老师)

1. cu 文件一般是用来写 cuda 的核函数
2. 在 .vscode/setting.json 中配置 *.cu : cuda-cpp，可以使得代码被正确解析
3. Makefile 中，cu 交给 nvcc 进行编译
4. cu 文件可以当作正常 cpp 写即可，它是 cpp 的超集，兼容支持 cpp 的所有特性
5. cu 文件中引入了一些新的符号和语法

• __global__ 标记，核函数标记
  • 调用方必须是 host
  • 返回值必须是 void
  • 例如：__global__ void kernel(const float* pdata, int ndata)
  • 核函数必须以 kernel<<<gridDim, blockDim, bytesSharedMemorySize, stream>>>(pdata, ndata) 的方式启动
  • 其参数类型是：<<<dim3 gridDim, dim3 blockDim, size_t bytesSharedMemorySize, cudaStream_t stream>>>
  • dim3 有默认构造函数 dim3(int x, int y=1, int z=1)
  • 因此当直接赋值为 int 时，实则定义了 dim.x = value, dim.y = 1, dim.z = 1
  • 其中 gridDim，blockDim，bytesSharedMemory，stream 是线程 layout 参数
  • 如果指定了 stream，则把核函数加入到 stream 中异步执行
  • pdata 和 data 则是核函数的函数调用参数
  • 函数调用参数必须传值，不能传引用等。参数可以是类类型等
    • 核函数的执行无论 stream 是否为 nullptr，都将是异步执行
  • 因此在核函数中进行 printf 操作，你必须进行等待，例如 cudaDeviceSynchronize 或者 cudaStreamSynchronize，否则你将无法看到打印的信息
• __device__ 标记，设备调用的函数
  • 调用方必须是 device
• __host__ 标记，主机调用的函数
  • 调用方必须是主机
• 也可以 __deivce__ __host__ 两个标记同时有，表明该函数可以设备也可以主机
• __constant__ 标记，定义常量内存
• __shared__ 标记，定义共享内存

6. 通过 cudaPeekAtLastError/cudaGetLastError 函数，可以捕获核函数是否出现错误或异常
7. 内存索引的计算公式

position = 0
for i in range(6):
   position *= dims[i]
   position += indexs[i]

8. buildin 变量，即内置变量，通过 ctrl+鼠标左键点进去查看定义位置

   • 所有核函数都可以访问，其取值由执行器维护和改变
   • gridDim[x, y, z]：网格维度，线程布局的大小，是核函数启动时指定的
   • blockDim[x, y, z]：块维度，线程布局的大小，是核函数启动时指定的
   • blockIdx[x, y, z]：块索引，对应最大值是 gridDim，由执行器根据当前执行的线程进行赋值，核函数内访问时已经被配置好
   • threadIdx[x, y, z]：线程索引，对应最大值是 blockDim，由执行器根据当前执行的线程进行赋值，核函数内访问时已经被配置好
   • Dim 是固定的，启动后不会改变，并且是 Idx 的最大值
   • 每个都具有 x、y、z 三个维度，分别以 z、y、x 为高低顺序

9. 关于 thread，grid，block 和 threadIdx 概念

• 首先，我们可以先不严谨地认为，GPU 相当于一个立方体，这个立方体有很多小方块如下图

• 每个小块都是一个 thread，为了方便讨论，我们只考虑 2D 的，如下图

• 我们关心的是某一个 thread 的位置，比如上图中的黄色方块
• 它在 2D 的位置是 (blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y) = (1, 0, 1, 1)
• 如果将这个 2D 展开成 1D，这个黄色 thread 的 1D 位置就是 13
• 计算方式如下图
• 但是一般情况，为了简化问题，我们只需要用到 threadIdx.x，blockIdx.x，blockDim.x 这三个量即可，所以计算 idx 的公式如下：
• int idx = threadIdx.x + blockIdx.x * blockDim.x; 其表示的含义是要求 thread 的 1D index，先得知道在第几个 block 里，再知道在这个 block 里得第几个 thread

总结

本次课程学习了核函数，它是一个在 GPU 上并行计算的函数，由 __global__ 符号进行修饰说明。核函数与普通的函数不同，在调用时需要使用 <<<>grid, block>> 来指定 kernel 要启动的线程数量，而每个线程都有唯一的线程号 thread ID 来标识，关于线程的全局索引计算可以根据杜老师的方法，采用左乘右加的方式进行记忆。

除此之外，我们还要对线程结构有一定的了解，一个 kernel 启动的所有线程被称为一个 grid，而一个 grid 里面又有很多的 block，一个 block 里面还包含有很多线程。grid 只是逻辑层，SM（流式处理器）才是执行的物理层，SM 的基本执行单元是 warp（线程束），每个 warp 包含 32 个线程。

最后我们写了一个简单核函数案例了解了核函数的定义和启动，并使用 threadIdx、blockIdx、blockDim 等信息来实现并行计算。
标识，关于线程的全局索引计算可以根据杜老师的方法，采用左乘右加的方式进行记忆。