CUDA 平台与编程模型
GPU 是硬件，CUDA 不是 GPU 的别名。CUDA 是 NVIDIA 提供的编程平台：包含编程模型、编译工具、驱动接口、运行时、调试分析工具和加速库。
1. CUDA 软件栈
text
application
├── CUDA libraries: cuBLAS / cuFFT / ...
├── CUDA Runtime API: cudaMalloc, cudaMemcpy, kernel launch
├── CUDA Driver API: contexts, modules, lower-level control
├── NVIDIA driver
└── NVIDIA GPU hardware
CUDA Toolkit: nvcc + headers + libraries + tools- Driver 负责操作系统与 GPU 的底层交互；
- Runtime API 提供更方便的 C/C++ 接口，通常在内部使用驱动能力；
- Toolkit 是开发工具集合，不等于显卡驱动；
- nvcc 协调 host C++ 与 device code 的编译；
- PTX 是面向 NVIDIA GPU 的虚拟指令表示，可由驱动进一步生成目标机器码；
- cuBLAS 等库 是经高度优化的现成算法，通常优先于自己重写通用算子。
安装 Toolkit 后能编译，不代表机器一定有可运行 CUDA 的 GPU；只有驱动也不代表拥有完整开发工具链。
2. Host 与 Device
text
CPU / host GPU / device
allocate and prepare data device memory
| cudaMemcpy H2D ------------> |
| launch kernel -------------> | many threads execute
| cudaMemcpy D2H <-------------|
consume resultKernel 启动对 Host 通常是异步的；一次同步的数据拷贝或显式同步可能让 Host 等待。准确行为要看 API、内存类型与 Stream。
3. 最小 Kernel
cpp
__global__ void add(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 threads = 256;
int blocks = (n + threads - 1) / threads;
add<<<blocks, threads>>>(a, b, c, n);边界判断不可省略，因为向上取整后的逻辑线程数可能大于 n。
4. 编程模型映射到硬件
text
CUDA programming model NVIDIA hardware
Grid
├── Block 0 -------------------------> resident on one SM
│ ├── Thread 0 ├── Warp scheduler
│ ├── ... ├── registers
│ └── Thread 255 └── shared memory + units
└── Block 1 -------------------------> an available SM
threads in a block -- groups of usually 32 --> warps一个 Block 在一次驻留期间位于一个 SM；一个 SM 可同时驻留多个 Block；不同 Block 的先后顺序通常不能假设。Block 是协作和同步边界，Warp 是硬件执行组织，二者不是同一层概念。
5. 同步边界
__syncthreads() 是 Block 内 barrier，并让同一 Block 的相关共享内存操作按规则可见。它不是全 Grid barrier。在条件分支中让部分线程到达 barrier、部分不到达，会造成错误。
跨 Block 的全局阶段常拆成多个 Kernel，因为前一个 Kernel 完成可作为下一阶段边界；更高级的协作机制有额外前提，不能默认使用。
深入原理与工程实践
前面的内容负责建立统一心智模型；下面把同一主题继续拆到执行过程、代码、性能代价与工程判断。
<!-- migrated-deep-dive:start -->
完整迁入：原 CUDA 编程模型全文
CUDA 编程模型——把矩阵乘法映射到 GPU / The CUDA Programming Model for Mapping Matrix Multiplication to GPUs
📅 创建时间：2026-06-02 🏷️ 标签：#CUDA #kernel #grid #block #thread #nvcc 📚 前置知识：[[03-gpu-architecture]]（GPU 架构） 📚 相关知识：[[05-cuda-kernel-and-memory]]（CUDA 内存） [[06-cuda-optimization]]（性能优化）
先抓住直觉
CUDA 的核心不是新语法，而是把一份工作拆给许多编号不同的线程。Host 负责准备和发号施令，Device 负责批量计算；每个线程根据自己的编号找到要处理的数据。
- 必须理解：Host/Device 分工；Grid、Block、Thread 的包含关系；线程索引如何映射到数据。
- 用到再查：函数修饰符全集、CMake 配置和错误码。
- 最小目标：先看懂向量加法，再进入矩阵乘法；不要一开始就手写高性能 GEMM。
场景：第一个 CUDA 程序——向量加法
┌─────────────────────────────────────────────────────────────┐
│ │
│ 你已经理解了 GPU 的架构： │
│ - 6912 个 CUDA 核心 │
│ - SM → Warp → Thread 的调度层次 │
│ - SIMT 执行模型 │
│ │
│ 现在，你要把这个简单的 C++ 代码改成 CUDA 版本： │
│ │
│ // C++ 版本 │
│ for (int i = 0; i < N; i++) { │
│ c[i] = a[i] + b[i]; │
│ } │
│ │
│ 目标：在 GPU 上并行执行这段计算 │
│ │
└─────────────────────────────────────────────────────────────┘第1节：Host 和 Device——CUDA 的核心概念
两种代码，两种运行环境
┌─────────────────────────────────────────────────────────────┐
│ CUDA 程序的两部分 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Host（C++ 代码，运行在 CPU 上）： │
│ - 数据准备：把数据从 CPU 内存拷贝到 GPU 显存 │
│ - 调用 kernel：告诉 GPU 要执行什么计算 │
│ - 结果回收：把结果从 GPU 显存拷回 CPU 内存 │
│ - 控制流：for/if 等逻辑 │
│ │
│ Device（CUDA 代码，运行在 GPU 上）： │
│ - 具体计算：每个线程执行一小部分计算 │
│ - 线程协作：同 SM 内线程共享数据 │
│ │
│ 形象理解： │
│ Host = 包工头（CPU）：分配任务、准备材料、收集结果 │
│ Device = 建筑队（GPU）：大规模并行施工 │
│ │
└─────────────────────────────────────────────────────────────┘代码分布示例
cpp
// ================================================================
// main.cpp（C++，Host 端）
// ================================================================
#include <cuda_runtime.h> // CUDA 运行时 API
#include <stdio.h>
int main() {
const int N = 1000000;
float *h_a, *h_b, *h_c; // h_ = host 内存（CPU）
float *d_a, *d_b, *d_c; // d_ = device 内存（GPU）
// Step 1: 分配 Host 内存
h_a = (float*)malloc(N * sizeof(float));
h_b = (float*)malloc(N * sizeof(float));
h_c = (float*)malloc(N * sizeof(float));
// Step 2: 初始化数据
for (int i = 0; i < N; i++) {
h_a[i] = i * 1.0f;
h_b[i] = i * 2.0f;
}
// Step 3: 分配 Device 内存（GPU 显存）
cudaMalloc(&d_a, N * sizeof(float));
cudaMalloc(&d_b, N * sizeof(float));
cudaMalloc(&d_c, N * sizeof(float));
// Step 4: CPU 数据拷贝到 GPU
cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);
// Step 5: 调用 kernel（启动 GPU 计算）
// 解释见下一节
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
// Step 6: GPU 结果拷贝回 CPU
cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);
// Step 7: 验证结果
for (int i = 0; i < 10; i++) {
printf("c[%d] = %f\n", i, h_c[i]);
}
// Step 8: 释放内存
free(h_a); free(h_b); free(h_c);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}cpp
// ================================================================
// vector_add.cuh（CUDA 头文件）
// ================================================================
#ifndef VECTOR_ADD_CUH
#define VECTOR_ADD_CUH
// __global__ 声明的函数从 CPU 调用，在 GPU 上执行
void add(float* a, float* b, float* c, int N);
#endifcpp
// ================================================================
// vector_add.cu（CUDA 源代码，Device 端）
// ================================================================
// CUDA 头文件
#include "vector_add.cuh"
#include <cuda_runtime.h>
// __global__: 从 CPU 调用，在 GPU 上运行的函数（称为 "kernel"）
// 这是 CUDA 编程的核心：一个 kernel 被 N 个线程并行执行
__global__
void add_kernel(float* a, float* b, float* c, int N) {
// 每个线程计算一个元素的结果
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 边界检查：防止数组越界
if (i < N) {
c[i] = a[i] + b[i];
}
}
// 包装函数（Host 端调用这个，它再调用 kernel）
void add(float* a, float* b, float* c, int N) {
// 计算 grid 和 block 的维度
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 启动 kernel：<<<blocksPerGrid, threadsPerBlock>>>
add_kernel<<<blocksPerGrid, threadsPerBlock>>>(a, b, c, N);
// 注意：kernel 启动是异步的！
// CPU 不会等待 GPU 完成就继续执行下一行
// 下一行立即执行，除非调用 cudaMemcpy 或 cudaDeviceSynchronize
}编译和运行
bash
### 编译
nvcc vector_add.cu main.cpp -o vector_add
### 运行（需要机器上有 NVIDIA GPU）
./vector_add
### 如果没有 GPU，可以模拟运行（需要 NVIDIA 开发驱动）
### CUDA Toolkit 自带 cuda-gdb 和 compute-sanitizer第2节：Grid / Block / Thread——CUDA 的三层线程组织
线程层次结构
┌─────────────────────────────────────────────────────────────┐
│ CUDA 线程层次结构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Grid（整个 GPU） │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ Block 0 Block 1 Block 2 Block 3 ... │ │
│ │ ┌─────┐ ┌─────┐ ┌─────┐ ┌─────┐ │ │
│ │ │ T0 │ │ T0 │ │ T0 │ │ T0 │ │ │
│ │ │ T1 │ │ T1 │ │ T1 │ │ T1 │ │ │
│ │ │ ... │ │ ... │ │ ... │ │ ... │ │ │
│ │ │ T255 │ │ T255 │ │ T255 │ │ T255 │ │ │
│ │ └─────┘ └─────┘ └─────┘ └─────┘ │ │
│ │ 256线程 256线程 256线程 256线程 │ │
│ │ │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ Grid 的维度：gridDim.x, gridDim.y, gridDim.z │
│ Block 的维度：blockDim.x, blockDim.y, blockDim.z │
│ │
└─────────────────────────────────────────────────────────────┘线程索引计算
cpp
// kernel 内部，每个线程通过内置变量知道自己负责哪个数据
// 内置变量速查：
// threadIdx.x/y/z 当前线程在线程块内的索引
// blockIdx.x/y/z 当前线程块在 grid 内的索引
// blockDim.x/y/z 每个线程块的维度
// gridDim.x/y/z grid 的维度
// 计算全局线程索引（唯一标识每个线程）
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 二维 Grid + 二维 Block 的例子：
// gridDim = (Gx, Gy), blockDim = (Bx, By)
// 全局索引 = (by * Bx + bx) * (Gy * Gx) + (ty * Bx + tx)
// 但实际代码中通常只用一维就够用了维度设置规则
┌─────────────────────────────────────────────────────────────┐
│ 维度设置规则 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Block 维度（每个 block 的线程数）： │
│ - 最大值：1024（硬件限制） │
│ - 常用值：128, 256, 512, 1024 │
│ - 推荐值：256（经验最优） │
│ │
│ Block 内线程数最好是 32 的倍数（Warp 大小） │
│ 因为 warp = 32 threads，必须整 warp 调度 │
│ 如果不是 32 的倍数，会浪费部分硬件 │
│ │
│ Grid 维度（block 的数量）： │
│ - 最大值：2^31 - 1（x方向）、65535（y）、65535（z） │
│ - Grid 大小 = block 数 × 每 block 线程数 = 总线程数 │
│ - 总线程数应该 >= GPU 核心数（通常远大于） │
│ │
│ 示例： │
│ A100 有 6912 个核心，N = 1000000 │
│ threadsPerBlock = 256 │
│ blocksPerGrid = ceil(1000000 / 256) = 3907 │
│ 总线程数 = 3907 × 256 = 1,000,192 > 6912 ✅ │
│ │
└─────────────────────────────────────────────────────────────┘第3节：CUDA 关键字——三种函数的区别
__global__ / __device__ / __host__
┌─────────────────────────────────────────────────────────────┐
│ CUDA 函数修饰符 │
├─────────────────────────────────────────────────────────────┤
│ │
│ __global__ void my_kernel(...) │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：CPU（Host） │ │
│ │ 执行位置：GPU（Device） │ │
│ │ 用途：kernel 入口（必须用这个声明） │ │
│ │ 注意：返回类型必须是 void │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ __device__ float my_device_func(float x) │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：GPU（Device） │ │
│ │ 执行位置：GPU（Device） │ │
│ │ 用途：GPU 内部调用的辅助函数 │ │
│ │ 注意：只能被 __global__ 或 __device__ 函数调用 │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
│ __host__ void my_host_func(...) ← 默认，省略不写 │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ 调用来源：CPU（Host） │ │
│ │ 执行位置：CPU（Host） │ │
│ │ 用途：普通 C++ 函数 │ │
│ │ 注意：可以同时有 __host__ 和 __device__（两边都编译）│ │
│ └───────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘变量修饰符
┌─────────────────────────────────────────────────────────────┐
│ CUDA 变量修饰符 │
├─────────────────────────────────────────────────────────────┤
│ │
│ __device__ ← 默认，全局内存，GPU 上 │
│ - 全局内存变量，所有线程可见 │
│ - 通过 cudaMalloc/cudaFree 管理 │
│ │
│ __shared__ ← 共享内存，同 SM 内线程 │
│ - 声明在 kernel 内部（或者用 shared 限定符） │
│ - 同一 block 内所有线程共享 │
│ - 速度极快（~1ns，接近寄存器） │
│ │
│ __constant__ ← 常量内存，只读 │
│ - 所有线程只读 │
│ - 编译时常量或运行时通过 cudaMemcpyToSymbol 设置 │
│ - 有硬件缓存（广播优化） │
│ │
│ __managed__ ← 统一内存（CUDA 6.0） │
│ - CPU 和 GPU 共享同一地址空间 │
│ - 自动在需要时迁移数据 │
│ │
└─────────────────────────────────────────────────────────────┘第4节：矩阵乘法的 CUDA 实现——从向量到矩阵
最简单的矩阵乘法
目标：C = A × B
A: M×K 矩阵
B: K×N 矩阵
C: M×N 矩阵cpp
// matrix_mul.cu
__global__
void matrix_mul_kernel(float* A, float* B, float* C,
int M, int N, int K) {
// 每个线程计算 C 的一个元素 C[row][col]
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < M && col < N) {
float sum = 0.0f;
for (int k = 0; k < K; k++) {
sum += A[row * K + k] * B[k * N + col];
}
C[row * N + col] = sum;
}
}cpp
// 调用方式（二维 Grid + 二维 Block）
void matrix_mul(float* A, float* B, float* C,
int M, int N, int K) {
// Block 维度：16×16 = 256 线程/block
dim3 threadsPerBlock(16, 16);
// Grid 维度：ceil(M/16) × ceil(N/16)
dim3 blocksPerGrid(
(N + threadsPerBlock.x - 1) / threadsPerBlock.x,
(M + threadsPerBlock.y - 1) / threadsPerBlock.y
);
matrix_mul_kernel<<<blocksPerGrid, threadsPerBlock>>>(
A, B, C, M, N, K);
}为什么一个线程算一个输出元素？
┌─────────────────────────────────────────────────────────────┐
│ 线程与输出的映射 │
├─────────────────────────────────────────────────────────────┤
│ │
│ C 是 M×N 的矩阵，有 M×N 个输出元素 │
│ 每个输出元素 C[i][j] 的计算是独立的！ │
│ │
│ ┌─────────────────────────────────────────────────────┐ │
│ │ Thread (block=0, thread=0) → 计算 C[0][0] │ │
│ │ Thread (block=0, thread=1) → 计算 C[0][1] │ │
│ │ Thread (block=0, thread=2) → 计算 C[0][2] │ │
│ │ ... │ │
│ │ Thread (block=1, thread=0) → 计算 C[1][0] │ │
│ │ ... │ │
│ │ Thread (block=M-1, thread=N-1) → 计算 C[M-1][N-1]│ │
│ └─────────────────────────────────────────────────────┘ │
│ │
│ M×N 个线程同时执行 = M×N 个输出同时计算 = 完美并行！ │
│ │
│ 这就是 GPU 编程的核心思想： │
│ 把计算拆分成 N 个独立的小任务，每个线程做一个 │
│ │
└─────────────────────────────────────────────────────────────┘更高效的共享内存版本（预告）
cpp
// 使用共享内存优化（大幅减少全局内存访问）
// 具体原理在第 5 章讲解，这里先看结构
__global__
void matrix_mul_shared_kernel(float* A, float* B, float* C,
int M, int N, int K) {
// 共享内存：同 block 内线程共享
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
float sum = 0.0f;
// 分块计算：每次加载一块到共享内存
for (int m = 0; m < (K + BLOCK_SIZE - 1) / BLOCK_SIZE; m++) {
// 线程协作加载到共享内存
if (row < M && (m * BLOCK_SIZE + threadIdx.x) < K)
As[threadIdx.y][threadIdx.x] = A[row * K + m * BLOCK_SIZE + threadIdx.x];
if (col < N && (m * BLOCK_SIZE + threadIdx.y) < K)
Bs[threadIdx.y][threadIdx.x] = B[(m * BLOCK_SIZE + threadIdx.y) * N + col];
// 等待所有线程加载完成
__syncthreads();
// 计算这一块
for (int k = 0; k < BLOCK_SIZE; k++)
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
// 等待所有线程计算完成，再加载下一块
__syncthreads();
}
if (row < M && col < N)
C[row * N + col] = sum;
}第5节：CUDA 错误处理——容易被忽略的坑
异步执行和错误检查
cpp
// ❌ 错误：cudaMalloc 的返回值不是错误码
float* d_a;
cudaError_t err = cudaMalloc(&d_a, N * sizeof(float));
if (err != cudaSuccess) {
printf("CUDA malloc failed: %s\n", cudaGetErrorString(err));
}
// ✅ 正确：kernel 启动是异步的！
// cudaGetLastError() 可以检查 kernel 启动是否有错
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
err = cudaGetLastError(); // kernel 启动后立即调用
if (err != cudaSuccess) {
printf("Kernel launch failed: %s\n", cudaGetErrorString(err));
}
// ✅ 正确：等待 GPU 完成后再检查结果
cudaDeviceSynchronize(); // 等待 GPU 所有任务完成
err = cudaGetLastError();
if (err != cudaSuccess) {
printf("Kernel execution failed: %s\n", cudaGetErrorString(err));
}完整错误检查封装
cpp
// 封装成宏，方便使用
#define CHECK_CUDA(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
printf("CUDA error at %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while (0)
// 使用
CHECK_CUDA(cudaMalloc(&d_a, N * sizeof(float)));
CHECK_CUDA(cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice));
add<<<blocks, threads>>>(d_a, d_b, d_c, N);
CHECK_CUDA(cudaGetLastError()); // kernel 启动后立即检查
CHECK_CUDA(cudaDeviceSynchronize()); // 等待完成
CHECK_CUDA(cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost));第6节：完整的项目结构示例
┌─────────────────────────────────────────────────────────────┐
│ 完整项目文件结构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ matrix_mul_project/ │
│ ├── CMakeLists.txt │
│ ├── include/ │
│ │ ├── matrix_mul.cuh ← kernel 声明 │
│ │ └── common.h ← 通用工具函数 │
│ ├── src/ │
│ │ ├── main.cpp ← Host 代码（C++） │
│ │ └── matrix_mul.cu ← Device 代码（CUDA） │
│ └── build/ │
│ └── (编译输出) │
│ │
└─────────────────────────────────────────────────────────────┘cmake
### CMakeLists.txt
cmake_minimum_required(VERSION 3.18)
project(matrix_mul LANGUAGES CXX CUDA)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_ARCHITECTURES 80) # A100
### CUDA Toolkit
find_package(CUDAToolkit REQUIRED)
### 头文件目录
include_directories(${CMAKE_SOURCE_DIR}/include)
### 可执行文件
add_executable(matrix_mul
src/main.cpp
src/matrix_mul.cu
)
target_link_libraries(matrix_mul PRIVATE CUDA::cudart)
### 启用 CUDA 扩展（CMake 自动处理 .cu 文件）
enable_language(CUDA)bash
### 编译
cd matrix_mul_project
mkdir build && cd build
cmake .. -DCMAKE_BUILD_TYPE=Release
cmake --build . --config Release
### 运行
./matrix_mul升华：CUDA 编程的核心心法
┌─────────────────────────────────────────────────────────────┐
│ CUDA 编程检查清单 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 1. 线程组织 │
│ ✅ 每个线程处理一个或多个数据元素 │
│ ✅ Block 维度设为 32 的倍数（warp 对齐） │
│ ✅ Grid 大小 >= GPU SM 数量（充分利用硬件） │
│ │
│ 2. 内存管理 │
│ ✅ cudaMalloc 分配 GPU 显存 │
│ ✅ cudaMemcpy 传输数据（注意方向 HostToDevice） │
│ ✅ cudaFree 释放显存 │
│ │
│ 3. 错误检查 │
│ ✅ 所有 CUDA 调用都要检查返回值 │
│ ✅ kernel 启动后立即检查 cudaGetLastError │
│ ✅ 验证结果前调用 cudaDeviceSynchronize │
│ │
│ 4. 性能意识 │
│ ✅ 线程数太少 → GPU 利用率低 │
│ ✅ 线程数太多 → 寄存器溢出，效率下降 │
│ ✅ 256 线程/block 是经验最优起点 │
│ │
└─────────────────────────────────────────────────────────────┘"AI 可查 vs 必须理解"清单
AI 可查：
✅ cudaMalloc/cudaMemcpy 的具体参数
✅ dim3 的具体用法
✅ CMake 中 CUDA 项目的具体配置
必须理解：
🔴 Host（CPU）负责调度，Device（GPU）负责计算
🔴 Grid / Block / Thread 的三层组织关系
🔴 线程索引计算：int i = blockIdx.x * blockDim.x + threadIdx.x
🔴 __global__ / __device__ / __host__ 的区别
🔴 kernel 启动是异步的，需要 cudaDeviceSynchronize 等待
🔴 为什么一个线程算一个输出元素 = 完美并行学习状态：🟡 开始学习
完整迁入：原并行计算版 CUDA 模型
CUDA 编程模型：Thread、Block、Grid 与内存协作 / CUDA Threads, Blocks, Grids, and Cooperative Memory Access
📅 创建时间：2026-07-20 🏷️ 标签：#CUDA #Kernel #ThreadBlock #GPU编程 📚 前置知识：[[05-gpu-architecture]] 📚 相关知识：[[/02-systems-and-performance/02-computer-architecture-and-hardware/04-cuda-programming-model]] [[/02-systems-and-performance/02-computer-architecture-and-hardware/05-cuda-kernel-and-memory]]
1. Host 与 Device
CUDA 程序通常同时运行在两个世界：
text
CPU / Host GPU / Device
准备数据 执行大量并行线程
分配显存 ───────────────→ 读取显存
启动 Kernel 完成计算
读取结果 ←─────────────── 写回显存CPU 负责控制流程，GPU 负责执行 Kernel。一次 GPU 加速是否值得，必须把数据传输和 Kernel 启动成本也算进去。
2. 第一个 Kernel
cpp
__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 threads = 256;
int blocks = (n + threads - 1) / threads;
vectorAdd<<<blocks, threads>>>(a, b, c, n);三个重要索引：
threadIdx：线程在线程块中的位置blockIdx：线程块在网格中的位置blockDim：每个线程块的尺寸
3. Grid、Block 与 Thread
text
Grid
├─ Block 0
│ ├─ Thread 0
│ ├─ Thread 1
│ └─ ...
├─ Block 1
└─ ...设计原则：
- Thread 处理最细粒度工作，例如一个元素
- Block 是能够共享内存并进行块内同步的线程组
- Grid 表示一次 Kernel 启动的全部线程块
不同 Block 之间默认不能通过普通 Barrier 同步，因此算法需要把跨 Block 协作放到后续 Kernel 或专门机制中。
4. 线程块怎样映射到硬件
线程块被调度到 SM，内部线程按 Warp 分组。常见一维 Block 大小为 128、256 或 512，但没有对所有 Kernel 都最优的固定值。
选择 Block 大小时要考虑：
- 是否为 Warp 大小的整数倍
- 每线程寄存器数量
- 每 Block Shared Memory 用量
- Kernel 的访存和计算特征
- 设备允许的最大线程数
Occupancy 高不等于性能一定高，它只表示活跃 Warp 相对于硬件上限的比例。
5. Shared Memory 协作
线程块可以先把数据从 Global Memory 搬到 Shared Memory，再重复使用：
cpp
__global__ void tiledKernel(const float* input, float* output) {
__shared__ float tile[256];
int global = blockIdx.x * blockDim.x + threadIdx.x;
tile[threadIdx.x] = input[global];
__syncthreads();
output[global] = computeWithNeighbors(tile, threadIdx.x);
}__syncthreads() 要求 Block 中所有仍然参与执行的线程到达同步点。把它放进线程条件不一致的分支可能导致错误或死锁。
6. 异步执行与 Stream
Kernel 启动通常对 CPU 是异步的。CUDA Stream 表示设备上的有序任务队列：
text
Stream 0：拷贝 A → Kernel A → 回传 A
Stream 1： 拷贝 B → Kernel B → 回传 B在硬件支持、数据独立且使用异步 API 的条件下，可以实现：
- 数据传输与计算重叠
- 多个 Kernel 并发
- 双缓冲流水线
Pinned Host Memory 有利于异步 DMA 传输，但锁页内存是有限资源，不应无限申请。
7. 错误检查
Kernel 错误可能在后续同步时才暴露。开发阶段至少检查：
cpp
kernel<<<blocks, threads>>>(...);
cudaError_t launchError = cudaGetLastError();
cudaError_t runtimeError = cudaDeviceSynchronize();生产代码应统一封装 CUDA API 错误处理，并记录设备、尺寸和调用上下文。
8. 常见初学者问题
- 忘记检查边界导致越界访问
- 频繁进行小块 Host-Device 拷贝
- 每个小操作都启动一个 Kernel
- 假设 Kernel 完成后 CPU 才继续执行
- 在分支中错误使用块内 Barrier
- 只关注线程数量，不关注访存模式
- 用默认 Stream 隐式同步整个流水线
核心总结
- CUDA 用 Grid、Block、Thread 描述大规模并行工作。
- Block 是共享内存与同步的基本协作范围。
- Kernel 启动、显存分配和数据传输都是成本。
- Stream 让传输与计算有机会重叠。
- 正确性、错误检查和边界处理先于性能优化。
下一篇：[[07-parallel-patterns]] <!-- migrated-deep-dive:end -->
面试速答
CUDA 与 GPU 的关系？ GPU 是硬件；CUDA 是 NVIDIA 针对其 GPU 的平台、模型、工具链和库生态。
Thread/Block/Grid 与 Warp/SM 的关系？ 前者是 CUDA 编程模型；Block 被调度驻留到 SM，其线程通常每 32 个组成 Warp 执行。
Runtime API 与 Driver API？ Runtime 更高层易用，Driver 更底层、控制更细；二者最终都依赖驱动与设备。
自测
- CUDA Toolkit、驱动和 Runtime 是否是同一个东西？
- 两个 Block 能否默认用
__syncthreads()同步？ - Kernel 启动后为什么不能立刻假定 Host 上结果可用？
答案：不是；不能；启动常异步，需要由拷贝、事件或显式同步建立完成关系。