CUDA 编程
CUDA 编程
1. CUDA 程序的基本结构
典型 CUDA 程序由两部分组成：
- Host 代码：运行在 CPU 上，负责分配内存、拷贝数据、启动 Kernel、同步结果。
- Device 代码：运行在 GPU 上，常写成
__global__Kernel 函数。
最小示例：
#include <cuda_runtime.h>
#include <cstdio>
__global__ void helloFromGpu() {
printf("Hello from block %d, thread %d\n", blockIdx.x, threadIdx.x);
}
int main() {
helloFromGpu<<<2, 4>>>();
cudaDeviceSynchronize();
return 0;
}启动语法：
kernel<<<gridDim, blockDim>>>(args...);含义是：启动 Grid，其中包含 gridDim 个 Block，每个 Block 包含 blockDim 个 Thread。
动手运行该示例时，将代码保存为 hello_cuda.cu，然后编译：
nvcc hello_cuda.cu -o hello_cuda运行：
./hello_cuda若系统找不到 nvcc，请先确认 CUDA Toolkit 已安装，且 nvcc 所在目录已加入 PATH。
2. Grid、Block、Thread
CUDA 的线程组织是层级结构：
Grid
└── Block
└── ThreadThread
Thread 是 CUDA 编程模型里最小的执行单位。每个 Thread 常处理一小份数据，例如数组中的单个元素。
常用内建变量：
threadIdx.x // 当前线程在 Block 内的 x 维编号
blockIdx.x // 当前 Block 在 Grid 内的 x 维编号
blockDim.x // 每个 Block 在 x 维有多少线程
gridDim.x // Grid 在 x 维有多少 Block一维数组常用全局线程编号：
int i = blockIdx.x * blockDim.x + threadIdx.x;Block
Block 是一组 Thread 的集合。Block 内线程具备以下特点：
- 通过 Shared Memory 共享数据。
- 使用
__syncthreads()做 Block 内同步。
- Block 调度到某个 SM 上执行。
- Block 不跨多个 SM。
Grid
Grid 是一次 Kernel Launch 创建的所有 Block 的集合。不同 Block 之间一般不能直接同步，也不能直接共享 Shared Memory。
需要跨 Block 同步时，常见做法是：
- 拆成多个 Kernel Launch，因为 Kernel Launch 之间天然有顺序关系。
- 使用 Cooperative Groups 等高级机制。
- 使用原子操作或全局内存协议，但要非常谨慎。
3. Warp：GPU 调度的基本执行单位
虽然 CUDA 编程时写的是 Thread，但 GPU 硬件实际调度的基本单位一般是 Warp。
在 NVIDIA GPU 上：
1 Warp = 32 Threads换言之，Block 中的线程按 Warp 切分。例如：
blockDim.x = 256
Warp 数量 = 256 / 32 = 8若 blockDim.x 不是 32 的倍数，末尾 Warp 中存在空闲 Lane。例如 100 个线程形成 4 个 Warp，其中末尾 Warp 只有 4 个有效线程。
Warp Divergence
同一 Warp 内的 32 个线程以 SIMT 方式执行。若它们走不同分支，将产生 Warp Divergence：
if (threadIdx.x % 2 == 0) {
// 一半线程执行该分支
} else {
// 另一半线程执行该分支
}硬件一般分别执行两个分支，同时屏蔽当前分支中不活跃的线程。因此分支越发散，执行效率越差。
4. SIMT：Single Instruction, Multiple Threads
SIMT 是 NVIDIA GPU 的核心执行模型，可理解为：
多个线程看起来各自独立执行，但硬件把同一 Warp 内的线程组织起来，使它们同时执行同一条指令。
这里的“同一条指令”不代表操作同一份数据。每个 Thread 都有自己的寄存器、自己的 threadIdx 和独立的执行上下文。同一 Warp 内的 32 个 Thread 可执行同一条访存或计算指令，但根据各自的线程编号访问不同地址。
例如：
int i = blockIdx.x * blockDim.x + threadIdx.x;
y[i] = x[i] * 2.0f;同一 Warp 内的 Thread 都执行 y[i] = x[i] * 2.0f 这条逻辑指令，但 threadIdx.x 不同，计算出的 i 也不同，因此访问的 x[i] 和 y[i] 通常不同。
SIMT 和 CPU 多线程不同：
- CPU 线程一般是重量级线程，有复杂调度和上下文切换。
- CUDA Thread 是轻量级线程，数量可非常多。
- GPU 依靠大量线程隐藏内存访问延迟。
SIMT 和 SIMD 也不同：
- SIMD 是一条指令显式操作多个数据 Lane，例如 AVX。
- SIMT 给程序员暴露的是 Thread 抽象，硬件内部把 Thread 组织成 Warp 执行。
5. SM：Streaming Multiprocessor
SM 是 NVIDIA GPU 的核心计算单元。单块 GPU 一般包含多个 SM。
可粗略理解为：
GPU
└── 多个 SM
├── Warp Scheduler
├── CUDA Cores
├── Register File
├── Shared Memory / L1 Cache
└── Load/Store Units, Special Function Units 等Block 和 SM 的关系
Block 调度到某个 SM 上执行，执行期间不迁移到其他 SM。
单个 SM 可同时驻留多个 Block，具体数量受以下因素限制：
- 每个 Block 的 Thread 数。
- 每个 Block 使用的 Register 数。
- 每个 Block 使用的 Shared Memory 大小。
- GPU 架构限制，例如单个 SM 最大 Thread 数、最大 Block 数。
这是 Occupancy 相关问题的来源。
Occupancy 是什么
Occupancy 指单个 SM 上实际驻留的活跃 Warp 数量，占该 SM 理论最大活跃 Warp 数量的比例。
例如，某个 SM 理论最多驻留 64 个活跃 Warp。当前 Kernel 由于 Register 或 Shared Memory 使用较多，只能驻留 32 个活跃 Warp，则：
Occupancy = 32 / 64 = 50%Occupancy 的意义在于隐藏访存延迟：某个 Warp 等待 Global Memory 数据返回时，SM 可切换执行其他活跃 Warp，从而减少执行单元空闲时间。
前面列出的 Thread 数、Register 数、Shared Memory 用量和 SM 架构上限，都影响 SM 能同时驻留多少活跃 Warp，进而影响 Occupancy。
注意，Occupancy 不是越高越好。若当前 Kernel 已有足够 Warp 隐藏访存延迟，继续提高 Occupancy 不一定提升性能；为提高 Occupancy 而压缩 Register 或 Shared Memory 使用，反而可能降低单线程效率或数据复用效果。
6. CUDA Core 与 Thread 的关系
- CUDA Thread 是编程模型里的逻辑执行单位。
- CUDA Core 是 SM 内部的硬件计算单元。
- Thread 数量一般远多于 CUDA Core 数量。
- Thread 不永久绑定 CUDA Core。
- Warp 中的指令调度到 SM 内的执行单元上执行。
可用类比理解：
CUDA Thread：任务
Warp：一组被一起调度的任务
SM：执行任务的车间
CUDA Core：车间里的具体计算工位因此，写 CUDA 程序时一般不需要直接关心“某个 Thread 在哪个 CUDA Core 上执行”。更重要的是合理组织：
- 每个 Block 有多少 Thread。
- 访存是否连续。
- 分支是否发散。
- Shared Memory 是否减少重复访问。
- Register 和 Shared Memory 是否限制 Occupancy。
7. Context：CUDA 上下文
CUDA Context 是 GPU 运行时的执行环境。它包含：
- 当前设备状态。
- 分配的 Device Memory。
- 已加载的 Kernel Module。
- Stream、Event 等运行时资源。
- 与当前 Host 线程相关的 CUDA 状态。
可将 Context 理解成 GPU 侧的“进程级运行环境”。
Runtime API 下的 Context
使用 CUDA Runtime API 时，Context 一般由运行时自动创建：
cudaSetDevice(0);
cudaMalloc(&ptr, size);首次触发设备操作时，CUDA Runtime 往往自动初始化对应 Device 的 Primary Context。
Driver API 下的 Context
使用 CUDA Driver API 时，则显式接触 Context：
cuInit(0);
cuDeviceGet(&device, 0);
cuCtxCreate(&context, 0, device);建议先使用 Runtime API，不必过早深入 Driver API。
8. Stream：并发执行队列
CUDA Stream 是 GPU 上的任务队列。同一 Stream 内的操作按提交顺序执行，不同 Stream 之间的操作可并发。
常见操作包括：
- Host 到 Device 的内存拷贝。
- Device 到 Host 的内存拷贝。
- Kernel Launch。
- Event 记录和等待。
默认 Stream 示例：
kernel<<<grid, block>>>(...);
cudaMemcpy(host, device, size, cudaMemcpyDeviceToHost);显式 Stream 示例：
cudaStream_t stream;
cudaStreamCreate(&stream);
kernel<<<grid, block, 0, stream>>>(...);
cudaMemcpyAsync(dst, src, size, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
cudaStreamDestroy(stream);Stream 的典型用途
Stream 常用于重叠数据拷贝和计算：
Stream 0: 拷贝 Chunk 0 -> 计算 Chunk 0 -> 拷回 Chunk 0
Stream 1: 拷贝 Chunk 1 -> 计算 Chunk 1 -> 拷回 Chunk 1若硬件支持 Copy Engine 和 Kernel 并发，可提升流水线效率。
要使 cudaMemcpyAsync 真正异步，Host 内存一般需要使用 Pinned Memory：
cudaMallocHost(&hostPtr, size);9. 内存层级
CUDA 程序性能主要取决于访存。
常见内存类型：
| 类型 | 位置 | 速度 | 作用域 | 典型用途 |
|---|---|---|---|---|
| Register | SM 内 | 最快 | 单个 Thread | 局部变量 |
| Shared Memory | SM 内 | 快 | 单个 Block | Block 内共享数据 |
| Global Memory | 显存 | 较慢 | 所有 Thread | 大数组、输入输出 |
| Constant Memory | 显存/缓存 | 快，适合广播 | 所有 Thread | 只读常量 |
| Texture Memory | 显存/缓存 | 针对空间局部性优化 | 所有 Thread | 图像、采样 |
最需要掌握的是：
- Global Memory 容量大但延迟高。
- Shared Memory 容量小但速度快。
- Register 最快，但过多 Register 将降低 Occupancy。
10. 示例一：向量加法
向量加法是经典 CUDA 例子：每个 Thread 处理单个元素。
#include <cuda_runtime.h>
#include <cstdio>
#include <vector>
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int main() {
const int n = 1 << 20;
const size_t bytes = n * sizeof(float);
std::vector<float> hA(n, 1.0f);
std::vector<float> hB(n, 2.0f);
std::vector<float> hC(n);
float *dA = nullptr, *dB = nullptr, *dC = nullptr;
cudaMalloc(&dA, bytes);
cudaMalloc(&dB, bytes);
cudaMalloc(&dC, bytes);
cudaMemcpy(dA, hA.data(), bytes, cudaMemcpyHostToDevice);
cudaMemcpy(dB, hB.data(), bytes, cudaMemcpyHostToDevice);
int block = 256;
int grid = (n + block - 1) / block;
vectorAdd<<<grid, block>>>(dA, dB, dC, n);
cudaMemcpy(hC.data(), dC, bytes, cudaMemcpyDeviceToHost);
printf("hC[0] = %f\n", hC[0]);
cudaFree(dA);
cudaFree(dB);
cudaFree(dC);
return 0;
}编译：
nvcc vector_add.cu -o vector_add运行：
./vector_add该例子体现最基本的 CUDA 工作流：
CPU 分配 Host 数据
CPU 分配 Device 数据
CPU 拷贝数据到 GPU
CPU 启动 Kernel
GPU 并行执行
CPU 拷贝结果回来11. 示例二：共享内存做归约
共享内存的经典场景是 Block 内协作。例如对数组求和，可由每个 Block 先算出局部和，再由后续 Kernel 或 CPU 合并。
下面是简化版 Block 内 Reduction：
#include <cuda_runtime.h>
#include <cstdio>
#include <vector>
__global__ void reduceSum(const float* input, float* partial, int n) {
// 每个 Block 分配一段 Shared Memory，供 Block 内所有 Thread 共同使用。
extern __shared__ float sdata[];
// tid 是 Thread 在当前 Block 内的编号，i 是该 Thread 对应的全局数组下标。
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
// 每个 Thread 从 Global Memory 读取一个元素到 Shared Memory。
// 超出数组范围的 Thread 写入 0，避免影响求和结果。
sdata[tid] = (i < n) ? input[i] : 0.0f;
// 等待 Block 内所有 Thread 完成数据加载。
__syncthreads();
// 折半归约：每轮把右半部分的值加到左半部分。
for (unsigned int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
if (tid < stride) {
sdata[tid] += sdata[tid + stride];
}
// 等待本轮累加完成，再进入下一轮。
__syncthreads();
}
// 归约完成后，sdata[0] 保存当前 Block 的局部和。
// 只需一个 Thread 写回结果，避免多个 Thread 重复写入。
if (tid == 0) {
partial[blockIdx.x] = sdata[0];
}
}
int main() {
const int n = 1 << 20;
const int block = 256;
const int grid = (n + block - 1) / block;
const size_t bytes = n * sizeof(float);
std::vector<float> hInput(n, 1.0f);
std::vector<float> hPartial(grid);
float *dInput = nullptr, *dPartial = nullptr;
cudaMalloc(&dInput, bytes);
cudaMalloc(&dPartial, grid * sizeof(float));
cudaMemcpy(dInput, hInput.data(), bytes, cudaMemcpyHostToDevice);
size_t sharedBytes = block * sizeof(float);
reduceSum<<<grid, block, sharedBytes>>>(dInput, dPartial, n);
cudaMemcpy(hPartial.data(), dPartial, grid * sizeof(float), cudaMemcpyDeviceToHost);
float sum = 0.0f;
for (float x : hPartial) {
sum += x;
}
printf("sum = %f\n", sum);
cudaFree(dInput);
cudaFree(dPartial);
return 0;
}关键点：
extern __shared__ float sdata[];表示动态 Shared Memory。
- Kernel 启动第三个参数
sharedBytes指定每个 Block 使用多少 Shared Memory。
__syncthreads()只同步同一 Block 内的 Thread。
- 最后每个 Block 输出 Partial Sum。
启动语法：
reduceSum<<<grid, block, sharedBytes>>>(dInput, dPartial, n);sharedBytes 表示每个 Block 分配的共享内存大小，不是整个 Grid 的总大小。
12. 示例三：共享内存优化矩阵乘法
矩阵乘法是 Shared Memory 的经典场景。
对于 C = A * B，每个 C[row][col] 都需要读取 A 的一行和 B 的一列。若每个 Thread 都直接从 Global Memory 读取，将产生大量重复访问。
使用 Tiled Shared Memory 的思路：
1. 每个 Block 负责 C 的 Tile。
2. Block 内线程协作把 A 的 Tile 和 B 的 Tile 加载到 Shared Memory。
3. 同步。
4. 在 Shared Memory 上做局部乘加。
5. 移动到下个 Tile。完整示例代码：
#include <cuda_runtime.h>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#define TILE 16
#define CUDA_CHECK(call) \
do { \
cudaError_t err = (call); \
if (err != cudaSuccess) { \
std::fprintf(stderr, "CUDA error %s:%d: %s\\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)
__global__ void matMulTiled(const float* A, const float* B, float* C, int N) {
// 每个 Block 复用 A 和 B 的一个 Tile。
// Shared Memory 位于 SM 内，比反复访问 Global Memory 更快。
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];
// 二维 Grid/Block 直接映射到矩阵坐标。
// blockIdx 决定当前 Block 位于结果矩阵 C 的哪个 Tile。
// threadIdx 决定当前 Thread 位于该 Tile 内的哪一行、哪一列。
// 两者组合得到全局坐标，当前 Thread 负责计算 C[row][col]。
int row = blockIdx.y * TILE + threadIdx.y;
int col = blockIdx.x * TILE + threadIdx.x;
// acc 保存当前 Thread 对 C[row][col] 的累加结果。
float acc = 0.0f;
// 沿矩阵乘法的公共维度 K 分块。
// 每轮加载 A 的一段列和 B 的对应一段行，累加到当前 C[row][col]。
for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
// 当前 Thread 负责加载 A[row][aCol] 和 B[bRow][col]。
int aCol = t * TILE + threadIdx.x;
int bRow = t * TILE + threadIdx.y;
// 越界位置填 0，处理 N 不能被 TILE 整除的情况。
As[threadIdx.y][threadIdx.x] = (row < N && aCol < N)
? A[row * N + aCol]
: 0.0f;
// B 的 Tile 也由 Block 内所有 Thread 协作加载。
Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N)
? B[bRow * N + col]
: 0.0f;
// 等待 A 和 B 的 Tile 全部加载到 Shared Memory。
__syncthreads();
// 使用 Shared Memory 中的 Tile 做局部乘加。
for (int k = 0; k < TILE; ++k) {
acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
// 等待本轮计算完成，再覆盖 Shared Memory 加载下个 Tile。
__syncthreads();
}
// 只写回矩阵范围内的元素。
if (row < N && col < N) {
C[row * N + col] = acc;
}
}
int main() {
const int N = 33;
const size_t bytes = N * N * sizeof(float);
std::vector<float> hA(N * N, 1.0f);
std::vector<float> hB(N * N, 2.0f);
std::vector<float> hC(N * N, 0.0f);
float *dA = nullptr, *dB = nullptr, *dC = nullptr;
CUDA_CHECK(cudaMalloc(&dA, bytes));
CUDA_CHECK(cudaMalloc(&dB, bytes));
CUDA_CHECK(cudaMalloc(&dC, bytes));
CUDA_CHECK(cudaMemcpy(dA, hA.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(dB, hB.data(), bytes, cudaMemcpyHostToDevice));
// 每个 Block 是 TILE x TILE 的二维线程布局。
dim3 block(TILE, TILE);
// Grid 也是二维布局，分别覆盖矩阵的列方向和行方向。
dim3 grid((N + TILE - 1) / TILE, (N + TILE - 1) / TILE);
matMulTiled<<<grid, block>>>(dA, dB, dC, N);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());
CUDA_CHECK(cudaMemcpy(hC.data(), dC, bytes, cudaMemcpyDeviceToHost));
const float expected = 2.0f * N;
bool ok = true;
for (float value : hC) {
if (std::fabs(value - expected) > 1e-5f) {
ok = false;
break;
}
}
std::printf("C[0][0] = %.1f, expected = %.1f\n", hC[0], expected);
std::printf("%s\n", ok ? "PASS" : "FAIL");
CUDA_CHECK(cudaFree(dA));
CUDA_CHECK(cudaFree(dB));
CUDA_CHECK(cudaFree(dC));
return ok ? 0 : 1;
}编译：
nvcc matmul_tiled.cu -o matmul_tiled运行：
./matmul_tiled预期输出类似：
C[0][0] = 66.0, expected = 66.0
PASS该例子体现 Shared Memory 的价值：
- 减少 Global Memory 重复读取。
- 使同一 Block 内的线程复用数据。
- 用更快的片上存储提升算术强度。
13. 选择 Block 大小的经验
常见入门选择：
128、256、512 Threads Per Block经验规则：
- 一般选择 32 的倍数，因为 Warp 包含 32 个线程。
- 256 是多数一维 Kernel 的稳妥起点。
- Block 太小可能无法充分利用 SM。
- Block 太大可能占用太多 Register 或 Shared Memory，反而降低并发驻留能力。
- 最终要结合 Profiling 工具判断。
常用 Profiling 工具：
nsys profile ./app
ncu ./app其中：
- Nsight Systems 适合查看 CPU/GPU 时间线、Stream 并发、拷贝和计算是否重叠。
- Nsight Compute 适合分析单个 Kernel 的访存、Occupancy、Warp Stall 等细节。
14. 错误检查
代码里建议加错误检查，否则 Kernel Launch 失败可能不容易发现。
示例宏：
#define CUDA_CHECK(call) \
do { \
cudaError_t err = (call); \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA error %s:%d: %s\\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
std::exit(EXIT_FAILURE); \
} \
} while (0)Kernel 后面检查：
vectorAdd<<<grid, block>>>(dA, dB, dC, n);
CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());注意：
cudaGetLastError()检查 Kernel Launch 是否成功。
cudaDeviceSynchronize()等待 GPU 执行完成，也能捕获异步执行期间的错误。
- 真实高性能代码中不一定每次都同步，但调试阶段有用。
15. 性能直觉
CUDA 性能优化一般围绕以下问题展开：
访存是否合并
连续 Thread 访问连续地址，一般更容易形成 Coalesced Memory Access。
好的模式：
int i = blockIdx.x * blockDim.x + threadIdx.x;
y[i] = x[i] * 2.0f;不好的模式：
y[i] = x[i * stride] * 2.0f;分支是否发散
同一 Warp 内线程尽量走相同路径。
是否有足够并行度
GPU 适合大量线程。只启动少量 Thread 往往无法发挥 GPU 性能。
Shared Memory 是否有效减少 Global Memory 访问
Shared Memory 不是“使用即加速”。若数据没有复用，或者引入过多同步，可能没有收益。
Kernel Launch 开销是否明显
大量小 Kernel 可能被 Launch Overhead 主导。可考虑合并 Kernel、使用 CUDA Graph，或调整任务粒度。
16. 常见概念关系总览
CUDA Context
└── 管理某个 GPU 上的运行环境、内存、Stream、Module 等资源
Stream
└── GPU 操作队列，同一 Stream 顺序执行，不同 Stream 可并发
Kernel Launch
└── 创建 Grid
└── Grid 包含多个 Block
└── Block 被调度到某个 SM
└── Block 包含多个 Thread
└── Thread 按 Warp 分组
└── Warp 一般包含 32 个 Thread
└── Warp 指令被 SM 内的 CUDA Core 等执行单元执行更简洁地说：
你写 Thread
CUDA 组织成 Block/Grid
硬件按 Warp 调度
SM 执行 Warp
CUDA Core 是 SM 里的计算资源17. 小结
- CPU 通过 CUDA Runtime/Driver API 控制 GPU。
- Context 是 GPU 的运行环境。
- Stream 是 GPU 操作队列。
- Kernel Launch 创建 Grid。
- Grid 包含 Block，Block 包含 Thread。
- Thread 在硬件上按 Warp 执行。
- Warp 被 SM 调度。
- CUDA Core 是 SM 内部执行算术指令的硬件资源，不等于 CUDA Thread。
- Shared Memory 是 Block 内线程协作和数据复用的重要工具。
大多数 CUDA 优化，本质上都是改善线程组织、访存模式、数据复用和硬件占用之间的平衡。