CUDA入门

要点速览（Abstract）：

  • 本文将涵盖一些CUDA的基础知识，笔者将会着重交流一下：分支分化及其避免方法、银行冲突及其避免方法。

一、CUDA概述

  1. CUDA 是什么？
  • NVIDIA推出的并行计算平台与编程模型
  • 扩展C/C++等语言支持GPU通用计算
  1. 为什么需要CUDA？
  • CPU与GPU架构差异（控制流 vs 数据流）
  • 并行计算加速场景：科学计算、深度学习、图像处理

二、CUDA编程模型的核心概念

  1. 线程层次结构

  • Thread（线程）：最小执行单位
    CUDA 中的线程执行相同的 Kernel 代码，但处理不同的数据元素，这种执行模式被称为"单指令多数据"（ SIMT ）。每个线程都有一个唯一的全局索引，用于确定它应该处理数据集中的哪个部分。

    每个 CUDA 线程都可以通过内置变量获取自己在整个线程层次结构中的位置。这些内置变量包括 threadIdx、blockIdx、blockDim 和 gridDim，它们分别表示当前线程在所属线程块中的索引、当前线程块在整个线程网格中的索引、线程块的维度以及线程网格的维度。这些变量都是预定义的结构体，通常包含 x、y（可能还有 z）三个分量，用于支持一维、二维或三维的线程组织方式。

    计算线程全局索引的公式为：global_idx = (blockIdx.x * blockDim.x) + threadIdx.x。这个公式假设使用一维的线程组织和索引计算。对于二维场景（如处理图像矩阵），可以使用 (blockIdx.y * blockDim.y + threadIdx.y) * width + (blockIdx.x * blockDim.x + threadIdx.x) 这样的公式将二维坐标转换为一维全局索引。

  • Block（线程块）：线程集合，共享内存与同步

    • 线程块的特性
      线程块是 CUDA 中线程的基本组织单元，一个线程块包含一组可以相互协作的线程。同一线程块内的线程可以通过共享内存和同步原语进行通信和协调，这是 CUDA 编程中实现复杂并行算法的基础。然而，每个线程块能够包含的线程数量是有限制的，这是由 GPU 硬件架构决定的。在大多数现代 GPU 架构中，一个线程块最多可以包含 1024 个线程，这个限制是在设计硬件调度器时确定的，开发者必须在这一约束内安排线程的组织方式。
      线程块在 GPU 上的执行是以线程束（Warp）为单位进行的。线程束是 GPU 硬件调度的基本单元，包含 32 个连续的线程。在一个线程束中，所有线程都执行相同的指令，但如果它们执行的条件分支不同，就会发生分支分化，导致部分线程等待其他线程完成，这种情况会降低执行效率。因此，设计 CUDA 程序时应当尽量避免同一个线程束内的分支分化，让所有线程执行相同的代码路径。理解线程束的概念对于优化 CUDA 程序的性能非常重要。
    • 线程块的同步与通信（后续会把__syncthreads()和__threadfence()放在一起讲）

  • Grid（线程网络）：Block集合，执行同一核函数
    线程网格是 CUDA 中最高层次的线程组织单位，它包含一个或多个线程块，所有线程块中的线程共同执行同一个 Kernel 函数。线程网格的维度由执行配置中的第一个参数指定，可以是一维、二维或三维。线程块在网格中的排列方式是确定的：对于一维网格，线程块按编号依次排列；对于二维网格，线程块按照先 x 方向后 y 方向的顺序排列。这种规则的排列方式使得线程块索引的计算变得简单直接。
    线程网格中的线程块是相互独立执行的，它们之间不能直接通信，也不能进行同步。GPU 的调度器会将线程块分配到不同的流式多处理器（ SM ）上执行，而线程块内部的线程则由同一个 SM 的线程调度器管理。这种设计简化了硬件设计，但同时也要求程序员在算法设计时考虑线程块之间的独立性。如果算法需要跨线程块的数据交换，必须通过全局内存作为中转，这在某些场景下可能会影响性能。

  • Kernel：内核函数
    Kernel 函数必须使用 global 限定符来声明，这是 CUDA 编译器识别的特殊标记。global 限定符表明这个函数可以从 CPU 端（主机端）调用，在 GPU 端（设备端）执行。Kernel 函数的返回类型必须是 void，因为从 CPU 的角度来看，Kernel 调用是异步的，CPU 不会等待 GPU 完成所有线程的执行就继续执行后续代码。这种设计允许 CPU 和 GPU 同时工作，最大化系统整体的计算效率。

    启动一个 Kernel 需要使用三重尖括号 <<< >>> 语法，这组尖括号被称为执行配置（Execution Configuration）。执行配置中需要指定线程网格和线程块的维度信息，这些参数决定了启动多少个线程以及如何组织这些线程。例如，<<<dimGrid, dimBlock>>> 表示启动一个包含 dimGrid 个线程块的网格，每个线程块包含 dimBlock 个线程。通过合理设置这些参数，可以控制并行计算的规模和粒度，以达到最佳的性能表现。

    Kernel 启动的第三个可选参数是一个指向共享内存的指针，用于在同一个线程块内的线程之间传递数据。这个参数在需要线程协作的场景中非常有用，例如规约操作（Reduction）或卷积计算中的窗口数据共享。Kernel 调用的实际语法形式为 kernel_name<<<grid, block, shared_mem_size, stream>>>(args)，其中 shared_mem_size 指定了为每个线程块分配的动态共享内存大小，stream 则指定了执行流，允许多个 Kernel 在同一个 GPU 上重叠执行。
    e.g:

__global__ void add_vectors(float* a, float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    float *a, *b, *c;
    // 内存分配和数据初始化代码...
    
    // 启动 Kernel：1 个线程块，1024 个线程
    add_vectors<<<1, 1024>>>(a, b, c, n);
    
    // 等待 GPU 完成计算
    cudaDeviceSynchronize();
    
    return 0;
}

几个重要的内存分配和初始化函数：

// 内存分配函数
cudaMalloc((void**)&d_A, sizeA); 
// 将数据从 Host 拷贝到 Device
cudaMemcpy(d_A, h_A, sizeA, cudaMemcpyHostToDevice);
// 配置线程块和网格
dim3 threadsPerBlock(16, 16);  // 16x16的线程块
dim3 blocksPerGrid((N + threadsPerBlock.x - 1) / threadsPerBlock.x, (M + threadsPerBlock.y - 1) / threadsPerBlock.y);

这段代码用于检查 kernel 启动错误，并在执行完成后同步，确保结果可用。

// 检查错误
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA Kernel Launch Error: %s\n", cudaGetErrorString(err));
    return EXIT_FAILURE;
}
cudaDeviceSynchronize();

三、CUDA基于硬件总览

  1. 流式多处理器（SM）
  • 流式多处理器（Streaming Multiprocessor，简称 SM）是 GPU 的基本计算单元，也是实现大规模并行计算的核心硬件模块。每个 SM 包含多个 CUDA 核心（也称为流式处理器或 SP）、共享内存、寄存器文件、常量缓存、纹理缓存以及指令调度单元等组件。
  • SM 的设计理念是"多线程并行执行"。每个 SM 能够同时管理并执行数百甚至数千个线程，这些线程被组织为线程束（Warp）进行调度。
  • SM 的资源分配策略决定了同时能在其上执行的线程块数量。每个 SM 拥有固定数量的寄存器文件、共享内存和其他资源，当用户启动 Kernel 时，CUDA 运行时需要根据每个线程块的资源需求计算一个 SM 上最多能容纳多少个线程块。这个过程称为" occupancy 计算"，它直接影响 GPU 的利用率和程序性能。如果 occupancy 过低，GPU 的计算能力将无法被充分利用；如果 occupancy 过高，每个线程可用的资源减少，可能导致寄存器溢出或共享内存不足，反而降低性能。
  1. 线程束（Warp）
  • 线程束（Warp）是 GPU 硬件调度的基本执行单元，由 32 个连续的线程组成。所有属于同一个线程束的线程在同一时刻执行完全相同的指令，这种执行模式被称为"单指令多线程"（SIMT）。
  • 当线程束中的线程执行条件分支语句（如 if-else）时，如果条件在某些线程上为真而另一些线程上为假，就会发生"分支分化"（Divergence）。在这种情况下，线程束需要分别执行两个分支的指令，某些线程可能需要等待其他线程完成其分支后才能继续执行，这会导致性能下降。因此，在编写 CUDA 代码时应当尽量避免同一个线程束内的分支分化，或者使用各种优化技术（如循环展开、谓词执行等）来减轻分支分化的影响。
  1. 高带宽内存（HBM）
  • 高带宽内存（High Bandwidth Memory，简称 HBM）是现代高端 GPU 采用的下一代内存技术，它通过堆叠式设计和宽总线接口提供了远超传统 GDDR6 内存的带宽。
  • HBM 的高带宽对于数据密集型应用（如深度学习训练）具有决定性意义。在大模型训练中，模型参数和中间激活值需要在内存和计算单元之间频繁传输。如果内存带宽不足，即使计算单元的理论算力再高，也会因为"数据饥饿"而无法发挥全部性能。HBM 技术通过提供超高带宽有效缓解了这一瓶颈，使得 GPU 能够持续地接收和处理数据，充分发挥 Tensor Core 的计算能力。
  1. 共享内存（SMEM）
  • 共享内存（Shared Memory，简称 SMEM）是位于 GPU 芯片上的高速内存，位于每个 SM 内部而非片外。共享内存的访问延迟通常只有几个时钟周期，远低于访问全局内存所需的数百个时钟周期。这种低延迟特性使得共享内存成为线程间通信和线程块内数据共享的首选机制。共享内存的容量是有限的，每个 SM 通常配备 48KB 到 128KB 的共享内存（具体大小取决于 GPU 架构和配置）。shared 限定符用于声明共享内存变量，这些变量在同一个线程块内的所有线程之间是共享的。共享内存的生命周期与线程块相同：当线程块开始执行时，共享内存被分配和初始化；当线程块执行结束时，共享内存被释放。
  • 具体用法：
    • 静态共享内存声明：
      • 静态共享内存是最简单的使用方式，直接在 Kernel 函数内部或外部使用 shared 限定符声明数组。例如：
__global__ void kernel_with_shared_memory(float* input, float* output, int n) {
    // 声明静态共享内存
    __shared__ float sdata[256];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 线程 0 将输入数据加载到共享内存
    if (idx < n) {
        sdata[tid] = input[idx];
    }
    
    // 同步确保所有线程完成数据加载
    __syncthreads();
    
    // 对共享内存中的数据进行规约操作
    // ... 规约代码 ...
    
    // 将结果写回全局内存
    if (tid == 0 && idx < n) {
        output[blockIdx.x] = sdata[0];
    }
}
  • 动态共享内存声明
    • 还要用extern修饰。指向未制定类型和大小的内存的指针。
// 动态共享内存版本
template<typename T>
__global__ void kernel_dynamic_shared(T* input, T* output, int n, size_t smem_size) {
    // 使用 extern 声明动态大小的共享内存
    extern __shared__ char sdata_char[];
    
    // 将 char 指针转换为需要的类型
    float* sdata = reinterpret_cast<float*>(sdata_char);
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 使用共享内存...
    if (idx < n) {
        sdata[tid] = input[idx];
    }
    
    __syncthreads();
    // ... 计算代码 ...
}

// 启动时指定共享内存大小
int shared_mem_size = 256 * sizeof(float);
kernel_dynamic_shared<<<grid, block, shared_mem_size>>>(input, output, n, shared_mem_size);
  • 共享内存的银行冲突（Bank Conflict）
    • 共享内存被组织为多个银行（Bank），每个银行可以独立提供服务。在现代 GPU 架构中，共享内存通常有 32 个银行，每个银行由连续的 4 字节地址组成。当多个线程同时访问同一个银行中的不同地址时，就会发生银行冲突（Bank Conflict），这些访问会被串行化执行，降低内存带宽利用率。避免银行冲突是优化 CUDA 程序的重要技巧之一。
    • 我来举个简单的例子，以下是一段矩阵转置的代码：
// 矩阵转置：读取行，写入列
// 读取时（无冲突）：
value = sharedArray[row * width + col];  // row相同，col连续

// 写入时（有冲突）：
sharedArray[col * height + row] = value;  // col相同，row连续

 冲突分析：
  假设width=32，height=32，一个warp处理一行：
    - 线程0: 写入地址 0*32 + 0 = 0 → bank 0
    - 线程1: 写入地址 1*32 + 0 = 32 → bank 0（冲突！）
    - 线程2: 写入地址 2*32 + 0 = 64 → bank 0（冲突！）
    - ...
    - 所有32个线程都访问bank 0 → 32-way冲突！
- 解决银行冲突的基本方法是确保同一线程束内的线程访问不同银行中的地址。最简单的策略是使用交错访问模式，例如通过 (tid & 31) * stride + (tid / 32) 而不是 tid 来索引共享内存数组。
-  为什么对呢？tid & 31 相当于对32取模从而得知线程在warp里的位置。tid/32可以得到线程所在的warp的索引。而stride我们一般取bank数量加一也就是33，33与32互质，于是随着tid&31遍历0到31，这个地址在32的模意义下也会跑遍完系，就能最大程度避免冲突。
  1. 寄存器（Register）
  • 寄存器是 GPU 中速度最快但容量最小的存储单元，位于每个 SM 的核心位置。每个 CUDA 核心拥有自己的寄存器文件，现代 GPU 的每个 SM 通常配备 64KB 到 256KB 的寄存器文件。
__global__ void kernel() {
   int local_var;           // 可能存储在寄存器
   float array[4];          // 小数组可能在寄存器
   // 寄存器特点：
   // 1. 线程私有，最快访问速度
   // 2. 数量有限（每线程典型256个）
   // 3. 生命周期=线程生命周期
}
  1. 线程组（WarpGroup）
  • 线程组（WarpGroup）是 Ampere 架构及以后 GPU 引入的一个新的线程组织概念，它将 4 个连续的线程束（128 个线程）组合在一起作为一个调度单元。
  • WarpGroup 使得同一线程组内的线程能够以更加协调的方式执行操作。传统的 CUDA 编程中，同一个线程束内的线程天然同步执行，但跨线程束的协调需要显式的同步原语。WarpGroup 提供了一种更灵活的机制，允许在更大的线程集合上进行协作计算，而不需要跨越线程束边界进行复杂的同步操作。这对于实现高效的矩阵运算和规约操作非常有用。
  • WarpGroup 与 Tensor Core 的协作是另一个重要特性。在使用 Tensor Core 进行矩阵运算时，操作通常以 4×4 或 8×8 的矩阵块为单位进行，这些块的大小正好与 WarpGroup 的规模相匹配。通过 WarpGroup，线程可以更高效地协调数据加载、计算和存储操作，充分发挥 Tensor Core 的峰值性能。在大模型推理中，WarpGroup 的异步复制能力可以显著提高内存带宽的利用率。
  1. 张量核心（Tensor Core）
  • 张量核心（Tensor Core）是 NVIDIA 从 Volta 架构开始引入的专门用于矩阵运算的硬件单元，它能够在一个时钟周期内完成一个完整的矩阵乘加运算（4×4 矩阵乘 4×4 矩阵加上另一个 4×4 矩阵）。与传统的 CUDA 核心相比，Tensor Core 的吞吐量提高了数倍甚至数十倍，这使得大规模矩阵运算的效率得到了质的飞跃。Tensor Core 的出现是深度学习在 GPU 上实现突破性加速的关键技术因素。
posted @ 2026-01-12 14:36  EternalEpic  阅读(24)  评论(0)    收藏  举报