02 CUDA 线程模型和显存模型

info

CUDA 是 NVIDIA 推出的并行计算平台，利用 GPU 强大的浮点运算能力加速 AI 与高性能计算任务。它提供线程模型、内存管理及编程接口，使开发者高效利用 GPU 进行大规模并行计算。我们在这一节课程中，同样会用到所学的共享显存知识来加速一个非常常见的 CUDA 核函数。

GPU 结构简介

我们来简单看一下 GPU 的结构，它是是围绕一个流式多处理器（SM）的扩展阵列搭建的，下图中包含以下的几个核心部件:

• LD/ST（Load/Store）单元:
  • 负责内存加载和存储操作，用于在寄存器和全局内存、共享内存之间传输数据。
• SFU（Special Function Unit）:
  • 专门执行特殊数学运算，比如三角函数、指数函数和倒数等，提升计算复杂函数的效率。
• CUDA Core:
  • 是 NVIDIA GPU 中的基本计算单元，负责执行基本的算术逻辑运算（如加法、乘法），是实现并行计算的核心组件。
• Dispatch Unit（分发单元）:
  • 负责将指令从线程束调度器发送到适当的执行单元（例如 CUDA Cores、LD/ST 或 SFU）。
• Warp Scheduler（线程束调度器）:
  • 负责管理和调度一个 SM（Streaming Multiprocessor）内的 warp（一组 32 个线程打包）；它选择准备好的 warp 并将它们的指令分派给相应的执行单元。

什么是 SM？

NVIDIA GPU 中的 SM（Streaming Multiprocessor，流式多处理器） 是 GPU 的核心执行单元。每个 SM 包含多个功能单元（如 CUDA 核心、内存访问单元等），并负责协调这些单元之间的协作，以实现高效的并行计算。在后续内容中我们将看到，一个线程块（thread block）一旦被分配到某个 SM 上，就会在其上完成全部执行过程，不会迁移到其他 SM，具体的调度机制将在本章后面的部分进行详细说明。

什么是 CUDA？

CUDA 出现的背景

• 在短短几年内，原本专为图形渲染而设计的可编程图形处理器（GPU），已经迅速演变为一种强大的并行数据计算设备。现代 GPU 的浮点运算能力已远远超过传统 CPU，在高性能计算领域展现出惊人的潜力。
• 这一转变背后的核心原因在于 GPU 架构的本质特征——它被专门优化用于高密度、大规模并行计算任务。最初，这种设计是为了满足复杂图形渲染中海量像素和顶点的实时处理需求；但随着硬件的发展和编程模型的进步，人们意识到 GPU 同样适用于其他具有高度并行特性的非图形计算任务，如和本门课强相关的 AI 加速方向，AI 计算任务通常涉及大量的浮点运算，其计算量大，数据频繁复用的特点正好契合了 GPU 的设计优势。
• 与 CPU 不同，GPU 更倾向于将芯片上的晶体管资源集中用于数据处理单元，而不是像 CPU 那样大量投入于缓存管理和复杂控制逻辑的设计。这种以计算为中心的架构使其在执行大规模并行任务时表现出极高的效率。

CUDA 技术栈

在以上的背景下，NVIDIA 顺势推出了一种 GPU 编程组件—— CUDA（Compute Unified Device Architecture） ，作为一种原生支持 GPU 编程的软硬件架构，CUDA 突破了过去必须通过图形 API 来使用 GPU 计算能力的限制，使得开发者可以直接在 GPU 上编写和执行通用计算程序。为了让大家对 Cuda 有一个直观理解，我们先来分层次看看它的结构

• 硬件驱动层 CUDA Driver:
  • 负责与 GPU 硬件通信，提供底层支持；
• 应用编程接口 API与运行时 Runtime:
  • 为开发者提供简洁易用的编程接口和执行环境；
• 高级数学库:
  • 提供优化的数学函数和算法，简化复杂计算任务的实现，例如:
  • CUBLAS - 用于线性代数运算的高性能库；
  • CUFFT - 用于快速傅里叶变换的库。

同时这里有一个重要的组件在图中没有体现，那就是 nvcc，也就是 NVIDIA CUDA 编译器（CUDA Compiler），它是 CUDA 工具链中的核心组件之一，专门用于编译含有 CUDA C/C++ 扩展语法的源代码文件（通常以 .cu 为扩展名）。

• CUDA C/C++ 是在标准 C/C++ 基础上进行扩展的一种语言，新增了用于编写设备端代码的关键字和语法结构，例如 global、device、host 等，这些扩展使得开发者可以在同一份代码中灵活地划分主机端与设备端的执行逻辑。
• CUDA C/C++ 还提供了丰富的运行时 API 和库函数（如 cudaMalloc、cudaMemcpy、cudaLaunchKernel 等），用于管理 GPU 内存、启动核函数以及实现 CPU 与 GPU 之间的数据传输与同步。

第一个 CUDA 程序

#include <cuda_runtime.h>
#include <iostream>

// =====================================================================
// 1. GPU kernel 函数
// __global__ 表示该函数运行在 GPU 上，由 CPU 端进行调用。
// <<<gridDim, blockDim>>> 会指定线程组织方式。
// =====================================================================
__global__ void hello_world(void) {
  // threadIdx.x 是 CUDA 内置变量，表示当前线程在 block 内的索引
  printf("thread idx: %d\n", threadIdx.x);

  // 只有 block 中的第 0 号线程执行这句
  if (threadIdx.x == 0) {
    printf("GPU: Hello world!\n");
  }
}

// =====================================================================
// 2. CPU 主函数（Host 代码）
// =====================================================================
int main(int argc, char **argv) {
  printf("CPU: Hello world!\n");

  // -------------------------------------------------------------------
  // 调用 GPU kernel
  // <<<1, 10>>> 表示：启动 1 个 block，每个 block 有 10 个线程
  // 也就是共有 10 个 GPU 线程执行 hello_world()
  // -------------------------------------------------------------------
  hello_world<<<1, 10>>>();

  // -------------------------------------------------------------------
  // cudaDeviceSynchronize()：阻塞 CPU，等待 GPU 执行完所有任务
  // 没有这句时，CPU 可能在 GPU 打印完成前就退出了程序。
  // -------------------------------------------------------------------
  cudaDeviceSynchronize();

  // -------------------------------------------------------------------
  // 检查 kernel 是否执行成功
  // cudaGetLastError() 会返回上一次 CUDA API 调用是否成功
  // -------------------------------------------------------------------
  if (cudaGetLastError() != cudaSuccess) {
    std::cerr << "CUDA error: "
              << cudaGetErrorString(cudaGetLastError())  // 错误原因文本
              << std::endl;
    return 1;
  } else {
    std::cout << "GPU: Hello world finished!" << std::endl;
  }

  std::cout << "CPU: Hello world finished!" << std::endl;
  return 0;
}

可以看到上述代码可以分解为两部分，Cuda 配套的编译器 nvcc 会自动识别并分离两部分代码，这二者的区别主要是主机端用于控制程序流程，管理资源(cudaMalloc 等)以及调用 GPU 设备端代码（核函数），而 GPU 代码则负责 GPU 上具体的并行计算任务。

1. CPU 主机端代码：通常由于 C/C++编写，它将由本地的 C/C++编译器（如 GCC、Clang 或 MSVC）进行编译.
2. GPU 设备端代码：使用 CUDA C/C++ 扩展语法 编写，由 nvcc 编译为可在 GPU 上执行的中间代码或 PTX 指令.

我们再来整体理解一下

• 在上面的代码中，我们看到一个特殊的函数调用：hello_world<<<1, 10>>>();
  • 这个语法用于启动名为 hello_world 的核函数
  • 并指定其执行配置为 1 个线程块（block），每个线程块中包含 10 个线程（thread）。
• 被 global 修饰的函数是运行在 GPU 上的设备端函数，在 CUDA 中被称为 kernel（核函数）。核函数是实现并行计算的核心机制，它将由多个线程同时执行，从而实现高度并发的计算任务。
• 在这个示例中，我们首次使用了一个重要的内置变量: threadIdx.x。
  • 该变量表示当前线程在其所属线程块中的一维索引值。
  • threadIdx 是一个三维结构体，包含 .x、.y .z 三个成员，分别对应线程在三个维度上的位置索引，便于在多维数据结构中进行寻址。
• 值得一提的是，主机（CPU）通常可以独立于设备（GPU）操作，并且不会阻塞自身执行。
  • 一旦核函数被启动，CUDA 运行时会立即返回控制权给主机端代码，从而实现主机与设备之间的异步执行。
  • 也正因如此，在主机代码中我们常常需要调用 cudaDeviceSynchronize() 函数，以显式等待之前启动的核函数执行完成，并确保可以正确获取执行结果或错误信息。

我们随后会看到，在 CUDA 中，Grid 里的所有线程都会执行同一个 kernel，所以如果 kernel 里有 printf，每个线程都会打印一次。

Grid（网格）
 └── Block（线程块）
      └── Thread（线程）

Grid
├── Thread Block
├── Thread Block
└── Thread Block
    ├── Thread1
    └── Thread2

# 例如 hello_world<<<1, 10>>>();
# 这里 1 就是 Grid 里有 1 个 Block
# 10 就是 每个 Block 里有 10 个 Thread

执行模型 - 线程分层和管理

执行模型的分层

为了方便编程，CUDA 提供了层次化的线程组织方式。每个线程都通过一个三维向量 threadIdx 来标识，允许我们以一维、二维或三维的方式组织线程，构成一个线程块（Thread Block）。这种方式使得开发者可以更自然地对向量、矩阵或三维数据进行并行计算。

• 图里的箭头代表 CUDA 的线程。线程组成线程块 (Thread Block)，多个线程块组成网格 (Grid)。Kernel 启动后，所有线程都属于这个 Grid。
• Grid 可以是一维、二维或三维的 Block 结构，Block 内的线程也能是多维的。但一个 Block 里的所有线程必须一起运行在同一个 SM (Streaming Multiprocessor) 上，因此线程数量受 SM 资源限制 (寄存器、共享内存等)。
• 当一个线程块 (Block) 被分配到某个 SM (Streaming Multiprocessor) 后，它会始终在该 SM 上执行，不会被迁移到其他 SM。多个 Block 可以同时分配到不同 SM 上，也可以多个 Block 共享同一个 SM，从而实现真正的并行。
• 在 Block 内，线程会按照 32 个线程为一组组合成线程束 (Warp)。Warp 中的 32 个线程同时执行同一条指令，但处理各自的数据。每个线程都有独立的程序计数器和寄存器。
• SM 会把一个 Block 切分成多个 Warp，再由 Warp Scheduler 选择可执行的 Warp，并把指令送到 Dispatch Unit，最终由硬件单元 (CUDA Cores、LD/ST、SFU 等) 来执行。

执行模型中软硬件的对应关系

SM = GPU 里的"小型执行器/小型计算引擎"

• 整个 GPU = 大工厂
• 每个 SM = 工厂里的一个“车间”
• 每个 SM 里有很多 CUDA Cores = 车间里的工人
• Warp = 一次派出去干活的 32 人小队
• Block = 一批任务
• Grid = 全部任务的集合

GPU
├── SM0
│    ├── 32-thread Warp
│    └── CUDA Cores
├── SM1
│    ├── Warp
│    └── CUDA Cores
└── SM2
     ├── Warp
     └── CUDA Cores

深入理解线程束 Warp

• 在线程块启动后，Block 内的所有线程会被划分成多个线程束 (warp)，每个 warp 固定包含 32 个线程。同一 warp 内的线程会同步执行同一条指令，但处理的数据各不相同，从而实现并行计算。
• 每个 SM 拥有固定数量的寄存器和共享内存。寄存器按线程分配，共享内存按线程块分配。当线程使用更多寄存器或线程块占用更多共享内存时，SM 能同时驻留的 Block 数会减少，因为硬件资源有限，从而影响可并行调度的 warp 数量。

• 在 线程束 (warp) 的执行过程中，可能会因为某些资源尚未就绪而发生 阻塞。例如，当一个 warp 需要从 全局显存 (global memory) 加载大量数据时，在数据返回之前，该 warp 会进入 等待状态。
• 由于 显存访问延迟远高于计算延迟，这种等待会显著降低执行效率。因此，GPU 的硬件设计将 隐藏访存延迟 (latency hiding) 作为提升性能的核心策略。其基本思想是：当某个 warp 因数据未准备好而阻塞时，SM 会立即 切换到另一个已经准备就绪的 warp 继续执行，避免硬件闲置，从而提升整体吞吐量。
• 一个 warp 若想在 SM 上被实际调度执行，必须满足以下两个条件：
  • 32 个 CUDA Core (执行单元) 可用
  • 执行所需的全部资源（寄存器、共享内存、操作数等）都已就绪

用例子理解 CUDA 的执行模型

• 我们以一个 向量加法 (vecAdd) 的示例来理解 CUDA 中线程的组织方式。线程块 (block) 是 CUDA 中的逻辑组织单位；但从 硬件执行角度 来看，线程的调度和访问最终仍是按照一维结构进行的。
• 下面考虑一个使用 __global__ 声明的核函数 vecAdd。它包含四个参数：输入数组 A、B，输出数组 C，以及数据总量 N。启动这个核函数时，我们可以按照三层结构组织线程：网格 (grid) → 线程块 (block) → 线程 (thread)。因为这是一个一维计算问题，我们通常采用 一维线程布局。每个线程都必须计算自己要处理的数据索引，这个索引可通过以下公式得到：

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

• 在 host 端，我们通过如下方式启动 kernel：

vecAdd<<<numBlocks, BLOCK_SIZE>>>(d_A, d_B, d_C, N);

启动配置 <<<numBlocks, BLOCK_SIZE>>> 表示我们总共启动了 numBlocks × BLOCK_SIZE 个线程。从代码可知，这个乘积刚好等于 N，意味着 每个数据元素对应正好一个线程。

• 在线程执行时，上面的索引公式：

idx = threadIdx.x + blockIdx.x * blockDim.x

会为每个线程计算其在整个数据中的 全局位置。举个例子：

• 若 BLOCK_SIZE = 256
• 则 numBlocks = 4
• 假设一个线程处于 blockIdx.x = 1 且 threadIdx.x = 1
• 那么它的全局索引为：

idx = 1 × 256 + 1 = 257

这表示该线程将负责处理 A[257] + B[257] 并将结果写入 C[257]。

• 通过这种 线程一维索引方式，我们可以精准地为每个线程分配待处理的数据元素，使得整个向量加法的执行既规则又高效。

```cpp
// vecadd_large.cu
// 向量加法的大规模 CUDA 示例，包含详细注释

#include <cstdio>

#define BLOCK_SIZE 256 // 每个 block 中包含的线程数量

// -----------------------------------------------------------------------------
// 核函数：vecAdd
// 每个线程计算一个元素：C[i] = A[i] + B[i]
// 参数：
//   A, B, C ：输入与输出数组（位于 GPU 显存中）
//   N       ：数组总长度
// -----------------------------------------------------------------------------
__global__ void vecAdd(int *A, int *B, int *C, int N) {
  // 计算当前线程对应的全局索引 i
  // threadIdx.x  -> 当前线程在 block 内的编号
  // blockIdx.x   -> 当前线程块的编号
  // blockDim.x   -> 每个 block 内的线程数（BLOCK_SIZE）
  int i = threadIdx.x + blockIdx.x * blockDim.x;

  // 越界检查，避免访问数组范围之外的数据
  if (i < N) {
    C[i] = A[i] + B[i];
  }
}

int main() {
  int N = 100000; // 要处理的数组大小（10 万个元素）
  size_t size = N * sizeof(int); // 数组占用的字节大小

  // ---------------------------------------------------------------------------
  // 在 host（CPU）端分配内存
  // ---------------------------------------------------------------------------
  int *A = (int *)malloc(size);
  int *B = (int *)malloc(size);
  int *C = (int *)malloc(size);

  // 初始化 A 和 B 数组的值
  for (int i = 0; i < N; i++) {
    A[i] = i;
    B[i] = i * 2;
  }

  // ---------------------------------------------------------------------------
  // 在 device（GPU）端分配内存
  // ---------------------------------------------------------------------------
  int *d_A, *d_B, *d_C;

  cudaMalloc((void **)&d_A, size);
  cudaMalloc((void **)&d_B, size);
  cudaMalloc((void **)&d_C, size);

  // ---------------------------------------------------------------------------
  // 将数据从 host 拷贝到 device
  // ---------------------------------------------------------------------------
  cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
  cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);

  // ---------------------------------------------------------------------------
  // 计算需要多少个 block 才能覆盖所有 N 个元素
  // numBlocks = ceil(N / BLOCK_SIZE)
  // ---------------------------------------------------------------------------
  int numBlocks = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;

  // ---------------------------------------------------------------------------
  // 启动核函数：<<<gridSize, blockSize>>>
  // gridSize  = numBlocks
  // blockSize = BLOCK_SIZE
  // ---------------------------------------------------------------------------
  vecAdd<<<numBlocks, BLOCK_SIZE>>>(d_A, d_B, d_C, N);

  // ---------------------------------------------------------------------------
  // 将计算结果拷贝回 host 端
  // ---------------------------------------------------------------------------
  cudaMemcpy(C, d_C, size, cudaMemcpyDeviceToHost);

  // ---------------------------------------------------------------------------
  // 验证结果是否正确
  // ---------------------------------------------------------------------------
  for (int i = 0; i < N; i++) {
    if (C[i] != A[i] + B[i]) {
      printf("Error at index %d: Expected %d, Got %d\n", i, A[i] + B[i], C[i]);
      break;
    }
  }
  printf("Vector addition completed successfully.\n");

  // ---------------------------------------------------------------------------
  // 释放 host 和 device 的内存
  // ---------------------------------------------------------------------------
  free(A);
  free(B);
  free(C);
  cudaFree(d_A);
  cudaFree(d_B);
  cudaFree(d_C);

  return 0;
}

存储模型 - 显存分层和管理

tip

显存原本是 "显示专用内存"（Video RAM），因为用于存储显卡渲染图像的数据，所以叫显存。 后来 GPU 做通用计算，但名字沿用至今。

为什么存储结构都是分层的

• 内存访问与管理是所有计算体系结构的核心要素。 在现代加速器（如 GPU）设计中，内存访问效率往往直接决定了程序的最终性能，因为许多实际应用都受限于数据加载与存储的速度，而不是算力本身。
• 大容量且高速的内存造价昂贵、制造难度高，且无法被大规模部署。 因此，硬件通常只能提供有限数量的高性能内存（例如寄存器、共享内存），而将更高容量的部分交给相对慢速的存储器（如 L2、全局显存等）。在这种物理限制下，分层式内存体系成为最可行的架构。
• CUDA 提供了一套 分层且透明的内存模型，统一了主机 (CPU) 与设备 (GPU) 的存储结构，并清晰展示了 GPU 的多级内存架构。通过利用这些不同级别的内存（寄存器、共享内存、全局内存、常量内存等），开发者可以 显式控制数据布局和访问策略，从而最大化吞吐率并最小化延迟。
• 因此，理解并合理使用 CUDA 的分层内存模型，是构建高性能 GPU 程序的基础。 只有充分利用不同层级存储器的延迟与带宽特性，才能真正发挥 GPU 的并行计算潜力。

CPU 上的内存分层结构

为了更好地理解内存层次结构的重要性，我们可以先从 CPU 上的经典内存体系 入手。现代计算系统的内存结构通常遵循一个核心原则：越靠近处理器的存储速度越快、但容量越小；越远离处理器的存储容量越大、但访问速度越慢。

在典型的 CPU 架构中，我们通常会看到如下由快到慢、由小到大的多级内存层次结构：

• 寄存器 (Registers)：
  • 位于处理器核心内部；
  • 访问速度最快（纳秒级甚至更快）；
  • 数量极为有限；
  • 用于存放指令执行时的临时数据。
• 高速缓存 (Cache)：
  • 包括 L1、L2、L3 缓存；
  • L1 最快但最小，L3 最慢但最大；
  • 用于减少 CPU 访问主存的延迟；
  • 构成现代处理器性能提升的关键手段。
• 主存 (Main Memory / DRAM)：
  • 容量远大于缓存；
  • 访问延迟明显更高（数十到数百纳秒）；
  • 是 CPU 运行程序和数据的主要载体。
• 持久化存储 (SSD / HDD)：
  • 容量巨大；
  • 延迟比内存高出成百上千倍；
  • 通常用于长期数据存储，不属于 CPU 直接访问的高速存储层级。

这种分层结构的根本原因在于：制造高速、大容量、低能耗的存储硬件在现实中无法同时做到，因此必须通过分层方式进行速度与容量的权衡。

GPU 上的显存分层结构

在 CUDA 编程中，寄存器 (Register)、共享显存 (Shared Memory) 和 全局显存 (Global Memory) 是最重要的三个内存层级，它们在容量、访问速度与作用范围上呈现出典型的分层特性：

• 容量分层：寄存器容量最小，共享显存次之，全局显存最大；
• 速度分层：寄存器最快，共享显存较快，全局显存最慢；
• 作用域分层：寄存器仅线程可见，共享显存在一个 Block 内可见，全局显存在整个 GPU 范围可见。

总结来说，三者在特性排序上分别为：

• 容量：寄存器 < 共享显存 < 全局显存
• 速度：寄存器 > 共享显存 > 全局显存
• 作用域：线程 < Block < GPU 全局

内存类型	存储位置 (Location)	作用域 (Scope)	可读写性	生命周期	访问速度	申请方式
寄存器	片上 (on-chip) 在芯片内部	单个线程 (thread)	可读写	线程执行期间	极快	自动
共享显存	片上 (on-chip) 在芯片内部	线程块 (Block) 内所有线程	可读写	Block 执行期间	快	显式声明
全局显存	片外 (off-chip) 不在芯片内部	所有线程 (全局可见)	可读写	主机端配置周期	最慢	主机端分配

下图更加清晰地展示了 GPU 中各级显存的层次关系

申请和传输数据到显存中

在 CUDA 中，主机 (Host) 与设备 (Device) 拥有独立内存，因此需要显式分配显存并手动完成数据传输。核函数只能访问设备端内存，因此设备内存的申请与数据拷贝是 CUDA 开发的基础步骤。

设备端申请显存

• cudaMalloc(void** devPtr, size_t count) 分配 count 字节的设备全局显存，并通过 devPtr 返回指针。
• cudaFree(void* devPtr) 释放先前分配的显存。

主机与设备之间的数据传输

• cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind) 用于主机与设备之间的数据拷贝，kind 指定方向

  • cudaMemcpyHostToDevice：主机 → 设备
  • cudaMemcpyDeviceToHost：设备 → 主机
  • 其他方向较少使用（HostToHost / DeviceToDevice）

由于核函数运行在设备端，无法直接访问主机内存，否则会产生未定义行为。因此所有在设备端使用的数据都必须先通过 cudaMemcpy 拷贝过去。

典型数据传输流程

1. 使用 cudaMalloc 分配设备显存；
2. 用 cudaMemcpyHostToDevice 将主机数据复制到设备；
3. 执行核函数处理数据；
4. 用 cudaMemcpyDeviceToHost 将结果从设备拷回主机。

以上步骤构成最基本的 CUDA 显存管理与数据传输模式，几乎出现在所有 CUDA 程序中。

#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

int main() {
  int size = 4 * sizeof(int);     // 要拷贝的字节数（4 个 int）

  int *h_a = (int *)malloc(size); // Host 端数组 h_a
  int *h_b = (int *)malloc(size); // Host 端数组 h_b（用于存回结果）
  int *d_a;                       // Device 端数组指针

  // 初始化主机端数据
  for (int i = 0; i < 4; i++) {
    h_a[i] = i;
  }

  // 在设备端分配显存
  cudaMalloc((void **)&d_a, size);

  // 将数据从 Host → Device
  cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);

  // 将数据从 Device → Host
  cudaMemcpy(h_b, d_a, size, cudaMemcpyDeviceToHost);

  // 打印结果，确认拷贝正确
  for (int i = 0; i < 4; ++i) {
    printf("h_b[%d] = %d\n", i, h_b[i]);
  }

  // 释放内存
  cudaFree(d_a);
  free(h_a);
  free(h_b);

  return 0;
}

该程序演示了 Host（主机）与 Device（设备）之间进行显式内存管理与数据传输的基本流程：

1. 在主机端分配内存 h_a 和 h_b，并将 h_a 初始化为 0 到 3；
2. 使用 cudaMalloc 在设备端分配内存 d_a；
3. 使用 cudaMemcpy（Host → Device）将 h_a 的数据复制到设备内存 d_a；
4. 再次使用 cudaMemcpy（Device → Host）将设备上的数据复制回 h_b；
5. 输出 h_b 的内容，并释放主机与设备的所有内存。程序正确运行后，h_b 将打印出 0 到 3，与 h_a 的初始值一致。

利用共享显存加速 CUDA Kernel

• 每个 SM 包含一小块低延迟的片上内存（Shared Memory），供同一线程块 (Block) 内的所有线程共享，是实现线程间协作的关键资源。
• 与全局显存 (Global Memory) 相比，共享显存的性能优势显著：访问延迟低 20–30 倍、带宽高约 10 倍。因此合理使用共享显存能显著提升核函数执行效率。
• 在 CUDA 中，共享显存可根据需要进行静态或动态分配。下面以“归约 (Reduction)”操作为例展示共享显存的重要性。我们先看一个未使用共享显存的版本，它直接从全局显存读取数据。

什么是 Reduction 函数?

Reduction 是一种将 大量数据 通过某种运算折叠成 一个结果 的 并行算法，在 GPU 中广泛用于 求和、最大值、统计 等关键操作，是 CUDA 编程中的基础 高性能并行模式。

• Reduction 的高频使用场景
  • 并行程序最常做的事就是 求和、最大值、统计，例如 数组求和、向量最大值、Softmax 最大值、反向传播的大量逐元素求和、物理模拟的全局能量统计、图像直方图、LLM Attention 中的 sum-exp。
  • 这些操作数据量巨大，是 GPU 工作负载中的常见核心步骤。
• 为什么 Reduction 必须优化
  • naive reduction 会频繁访问 全局显存、重复读写数据、产生 同步开销，且 warp 容易出现 分支 (divergence)，从而形成显著 性能瓶颈。
  • 因此 CUDA 使用 共享显存缓存数据、树形归约 (tree-based reduction)、warp 内归约、避免 bank conflict 和减少分支 等技巧来加速，使 Reduction 充分发挥 GPU 的并行优势。

不使用共享显存的 Reduction 示例分析

• 在示例代码（course1/reduce_smem.cu）中，线程直接从全局显存 g_idata 加载输入数据，而不是将其缓存到共享显存中。
• 以 threadIdx.x = 1 的线程为例
  • 第一次读取索引 1；
  • 后续归约中读取 1 + 512 = 513；
  • 再下一轮读取 1 + 256 = 257；
  • 多轮迭代后，最终将该线程块的部分归约结果写入全局显存 g_odata。
• 整个归约过程中，线程始终通过全局显存进行数据访问，没有借助共享显存做缓存或数据重用。此处无需完全理解算法逻辑，关键在于理解：直接访问全局显存的延迟高且代价大，而共享显存可显著降低访问成本。

// 不使用共享显存版本
#include <iostream>

__global__ void reduceGmem(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;   // 当前线程在 block 内的索引

  // 越界保护：如果线程编号超出 n，则直接退出
  if (tid >= n)
    return;

  // 指向当前 block 对应的全局显存片段
  // 每个 block 处理 blockDim.x 个数据
  int *idata = g_idata + blockIdx.x * blockDim.x;

  // ---------------------
  // 以下是逐级归约 (Reduction)
  // 每一轮把远端的数据加到靠前位置
  // 所有操作都在全局显存 idata 上进行（没有使用共享显存）
  // ---------------------

  // 第 1 轮：把 tid+512 的数据加到 tid 上
  if (blockDim.x >= 1024 && tid < 512)
    idata[tid] += idata[tid + 512];
  __syncthreads(); // 保证所有线程完成再进入下一轮

  // 第 2 轮：tid + 256
  if (blockDim.x >= 512 && tid < 256)
    idata[tid] += idata[tid + 256];
  __syncthreads();

  // 第 3 轮：tid + 128
  if (blockDim.x >= 256 && tid < 128)
    idata[tid] += idata[tid + 128];
  __syncthreads();

  // 第 4 轮：tid + 64
  if (blockDim.x >= 128 && tid < 64)
    idata[tid] += idata[tid + 64];
  __syncthreads();

  // ---------------------
  // warp 内归约（tid < 32）
  // 注意 warp 内不需要 __syncthreads()
  // 使用 volatile 强制内存访问顺序，避免编译器优化导致错误
  // ---------------------
  if (tid < 32) {
    volatile int *vmem = idata;

    vmem[tid] += vmem[tid + 32];
    vmem[tid] += vmem[tid + 16];
    vmem[tid] += vmem[tid + 8];
    vmem[tid] += vmem[tid + 4];
    vmem[tid] += vmem[tid + 2];
    vmem[tid] += vmem[tid + 1];
  }

  // ---------------------
  // 写出每个 block 的归约结果
  // block 的最终结果存在 idata[0]
  // ---------------------
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

使用共享显存优化的 Reduction 示例分析

下面的代码使用了 共享显存 (Shared Memory) 对输入数据进行缓存，避免同一个线程多次从 全局显存 (Global Memory) 读取相同数据，借助共享显存的高带宽和低延迟显著提升归约效率。

__global__ void reduceSmem(int *g_idata, int *g_odata, unsigned int n) {
  // 声明共享显存空间（与线程块大小对应）
  __shared__ int smem[256];

  unsigned int tid = threadIdx.x;   // 当前线程在 block 内的索引

  // 越界保护：若线程编号 >= n，则不参与归约
  if (tid >= n)
    return;

  // 指向当前 block 对应的全局显存区段
  // 每个 block 处理 blockDim.x 个元素
  int *idata = g_idata + blockIdx.x * blockDim.x;

  // 将全局显存的数据加载到共享显存中
  // 提高后续归约阶段的数据访问速度
  smem[tid] = idata[tid];

  // 同步线程，确保所有共享显存数据已加载完毕
  __syncthreads();
}

• 本例中新增了 __shared__ 声明，线程块内的所有线程协作，将全局显存中的数据统一加载到共享显存中，实现高效的数据预取。
• __syncthreads() 用于 同步线程块内所有线程，确保共享显存中的数据已完全就绪，再进入下一步计算，避免数据竞争。
• 共享显存大小定义为 256，是因为主机端启动 kernel 时设置 blockDim.x = 256，线程数与共享显存槽位一一对应，这 256 个线程将自己的数据加载到共享显存中，随后所有线程可快速访问这些缓存数据。

Host 函数中的调用与计时

这个 main 函数主要完成三件事：

1. 在 Host 上准备输入数据：构造一个长度为 102400、全部为 1 的数组，方便验证归约结果是否正确。
2. 在 Device 上分配显存并拷贝数据：使用 cudaMalloc 为 d_idata1、d_idata2 分配显存，并通过 cudaMemcpy 将 Host 端的全 1 数组拷贝到 GPU 上。
3. 使用 CUDA Event 测量核函数执行时间：通过 cudaEventRecord、cudaEventSynchronize 和 cudaEventElapsedTime 精确统计 reduceSmem 的运行时间。

#include <iostream>
#include <cuda_runtime.h>

int main() {
  const int n = 102400;      // 输入数据规模
  int h_idata[n];            // Host 端输入数据

  // 1. 在 Host 上初始化输入数据：全 1 数组
  for (int i = 0; i < n; i++)
    h_idata[i] = 1;

  int *d_odata1, *d_odata2, *d_idata1, *d_idata2;

  // 2. 在 Device 上申请显存
  cudaMalloc((void **)&d_idata1, n * sizeof(int));
  cudaMalloc((void **)&d_idata2, n * sizeof(int));
  cudaMalloc((void **)&d_odata1, sizeof(int));  // 存放归约结果
  cudaMalloc((void **)&d_odata2, sizeof(int));

  // 3. 将 Host 数据拷贝到 Device（两个输入缓冲相同，便于对比不同 kernel）
  cudaMemcpy(d_idata1, h_idata, n * sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_idata2, h_idata, n * sizeof(int), cudaMemcpyHostToDevice);

  // 假设在别处定义好了 blockSize 和 numBlocks：
  // int blockSize = 256;
  // int numBlocks = (n + blockSize - 1) / blockSize;

  // 4. 创建 CUDA Event，用于测量 kernel 执行时间
  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  // 5. 记录起始事件，调用 reduceSmem 核函数
  cudaEventRecord(start);
  reduceSmem<<<numBlocks, blockSize>>>(d_idata2, d_odata2, n);
  cudaEventRecord(stop);

  // 6. 等待 GPU 完成 stop 事件之前的所有工作
  cudaEventSynchronize(stop);

  // 7. 计算起始和结束事件之间的时间差（单位：毫秒）
  float millisecondsSmem = 0.0f;
  cudaEventElapsedTime(&millisecondsSmem, start, stop);
  std::cout << "Time for reduceSmem: " << millisecondsSmem << " ms" << std::endl;

  // （实际代码中建议补充结果检查和资源释放）

  return 0;
}

结果比较

我们在执行了reduce_smem.cu之后可以得到两个核函数的时间对比，可以明显看出使用共享显存的版本在执行结果一致的情况下对计算时间有一定的提升

Time for reduceGmem: 0.189984 ms
Time for reduceSmem: 0.132096 ms
Sum1: 512
Sum2: 512

NVIDIA Nsight Compute

Nsight Compute 是 NVIDIA 的 CUDA Kernel 性能分析器，用来查看一个 Kernel 在 GPU 上到底慢在哪里、瓶颈是什么。

• 访存效率：global memory 是否合并访问、带宽是否达到上限
• Warp 行为：是否发生 branch divergence
• Cache 与 Shared Memory 使用情况
• 寄存器 / SM 占用率（occupancy）
• Kernel 性能类型：是 memory bound 还是 compute bound

Nsight Compute 的分析结果表明，使用共享显存（reduceSmem）后:

• 执行时间提升约 18%
• SM 计算吞吐率提升 272%
• 内存吞吐率提升 49%
• L1 缓存命中率显著提高，L2/DRAM 访问减少
• Warp 阻塞明显降低（特别是 global memory 相关阻塞）

主要瓶颈依然是:

• Long Scoreboard Stalls（等待全局显存加载）
• Barrier Stalls（线程同步带来的阻塞）

这说明共享显存带来了显著优化，但进一步性能提升需要:

• Warp-level reduction（减少 barrier）
• 更好的 memory coalescing
• 更少的全局显存访问