cuda编码入门学习笔记

          在日常深度学习和科学计算中，使用图形处理器(GPU)进行加速是一个常见的做法。CUDA (Compute Unified Device Architecture) 是英伟达公司提供的用于GPU编程的平台和编程模型。同时它是一种并行计算模型，允许开发人员使用标准C语言对GPU进行编程,也就是对c语言的一种扩充写法。CUDA的核心思想是将任务分解为多个线程，并在GPU上同时执行这些线程。GPU由多个处理器和全局内存组成，每个处理器可以执行多个线程，同时访问全局内存。通过并行执行多个线程，可以大大提高计算速度。学习参考链接：https://docs.nvidia.com/cuda/cuda-c-programming-guide/

（1）CUDA编程模型

     》 学习CUDA编程模型的基础知识：

• Grids（网格）：

  • 网格是 CUDA 编程中的最高级别的组织单位。它是一个三维的结构，用于管理并行执行的线程块（blocks）。
  • 网格由一个或多个线程块组成，并且可以在三个维度上指定大小（例如 (x, y, z)）。
  • 网格的大小决定了可以并行执行的线程块的数量。

• warp（线程束）：

  • Warp是一个执行单元的最小集合。在当前的NVIDIA GPU架构中，一个warp包含32个线程。
  • 在一个warp中的所有线程会同步执行相同的指令，即SIMT（Single Instruction, Multiple Threads）模型。简而言之，warp中的所有线程同时执行相同的指令，只是它们处理的数据可能不同。

• Blocks（线程块）：

  • 线程块是网格中的次级单位，它负责管理一组并行执行的线程。
  • 线程块是一个三维的结构，可以在三个维度上指定大小（例如 (x, y, z)）。
  • 每个线程块中的线程可以协作并共享内存，通常被设计为处理一块数据或执行一个子任务。
  • Block是线程的一个更大的组织单位。一个block包含多个warp，具体数量取决于block中的线程数。

• Threads（线程）：

  • 线程是 CUDA 编程中的最小执行单位，被组织成线程块中的集合。
  • 线程通常是一维的，它们可以通过特定的线程索引访问数据并执行操作。
  • 线程可以利用GPU的并行性，以高效地执行计算任务。

• Kernel 函数：

  • Kernel 函数是在GPU上并行执行的函数，它由CPU发起并在GPU上执行。
  • Kernel 函数由关键字 __global__ 声明，用来标识它们可以被并行执行。
  • 当一个Kernel函数被调用时，它在GPU设备上的多个线程中并行执行，每个线程执行同一段代码但处理不同的数据。
  • Kernel 函数通常用于执行大规模数据并行计算任务，如矩阵运算、图像处理等。
• dim3数据格式

在CUDA编程中，dim3通常用于指定CUDA内核的执行配置，包括网格的维度（gridDim）和每个线程块的维度（blockDim）。dim3提供了一种方便的方式来处理一维、二维或三维的并行计算任务。

dim3 dimBlock(256); // 一维线程块，大小为256，等价于256×1×1
dim3 blockSize(16, 16); // 二维线程块大小为16x16，等级于16×16×1
dim3 dimGrid(10, 20, 30); // 三维网格，大小分别为10, 20, 30

在内核启动时，如果某个维度的大小为1，CUDA运行时会理解为一维情况。例如：

kernel<<<gridSize, blockSize, 0>>>(args); // 内核启动调用

在这个调用中，即使gridSize和blockSize是dim3类型，如果它们被定义为一维或二维，CUDA运行时也会正确处理它们。

---

补充：在CUDA编程中，__global__和__device__,__host__以及__shared__是几个重要的关键字，用于标识函数在GPU上执行的不同方式。

• __device__ 用于声明在设备上（即GPU上）全局可见的变量或函数，__device__变量在所有线程中都是可见的，但它们存储在全局内存中，访问速度相对较慢。__device__函数可以在设备代码中被其他内核调用，类似于普通的C/C++函数。但是CUDA中用来标识在GPU上执行，但只能被设备调用的函数。__device__ 函数可以有返回值，并且可以被其他 __device__ 或 __global__ 函数调用。与 __global__ 函数不同，__device__ 函数不支持被主机（CPU）直接调用，因为它们是专门为在GPU设备上运行而设计的。
• __shared__用于声明在同一个线程块内所有线程共享的变量。其中__shared__变量存储在共享内存中，访问速度比全局内存快得多，适合于线程间的数据共享和同步。但是共享内存的大小有限，且在内核启动时分配，因此需要谨慎使用以避免超出内存限制。
• __global__用于声明CUDA内核函数，这些函数可以在主机代码(cpu)中调用，并在设备上执行。__global__函数通常用于执行大规模的并行计算任务，它们会被映射到多个线程和线程块上。但是其返回类型必须是void，想要拿到处理结果需要将其拷贝到host端才行。
• __host__用于声明函数在CPU上执行，是CUDA编程中对函数执行位置进行明确标注的一种方式。大多数情况下在主机上执行的函数默认不需要显式声明__host__修饰符，但有时为了代码的可读性和明确性，或者在组合使用多个修饰符时，会显式使用__host__。

---

在 CUDA 编程中，核函数（kernel function）是在 GPU 上执行的函数。了解如何定义和使用核函数的参数对于有效利用 GPU 资源至关重要。

》核函数定义和参数

• 核函数定义： 核函数使用 __global__ 修饰符来定义。核函数是从主机代码（CPU 上执行的代码）调用的，但实际在设备代码（GPU 上执行的代码）上运行。

__global__ void kernelFunction(parameters) {
    // 核函数代码
}

• 参数类型： 核函数的参数可以是基本数据类型、指针或 CUDA 内置类型。

__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements) {
    // 核函数代码
}

• 线程索引

在核函数中，threadIdx、blockIdx、blockDim 和 gridDim 是 CUDA 内置变量，用于获取线程和块的索引信息。

int i = blockDim.x * blockIdx.x + threadIdx.x;

》核函数调用和配置

• 核函数调用： 核函数在主机代码中使用特殊的语法调用，基本调用方式（两个值）称为网格（grid）和块（block）配置

kernelFunction<<<blocksPerGrid, threadsPerBlock>>>(parameters);

其中，blocksPerGrid为每个网格中块的数量，threadsPerBlock为每个块中线程的数量。其取值尽量为选择能够被32整除的数，以便最大限度地利用GPU的线程束（warp）。

常见的还有带共享内存的调用方式（三个值），其中sharedMemSize：为每个线程块分配的共享内存大小（以字节为单位）。

kernel<<<gridSize, blockSize, sharedMemSize>>>(args);

以及带执行流的调用方式（四个值），其中stream：指定的CUDA流（cudaStream_t类型），用于控制命令的异步执行。

ker