02-GPU编程模型¶
重要性：⭐⭐⭐⭐⭐ 实用度：⭐⭐⭐⭐⭐ 学习时间：2天 必须掌握：是
为什么学这一章？¶
在了解GPU硬件架构之后，我们需要理解如何在软件层面组织和调度GPU上的并行计算。GPU编程模型定义了线程的层次结构、内存访问方式和同步机制，是编写高效CUDA程序的理论基础。
学完这一章，你将能够： - ✅ 理解CUDA的线程层次（Grid → Block → Thread） - ✅ 掌握线程索引计算方法 - ✅ 理解Warp执行模型和分支分歧 - ✅ 编写基本的CUDA核函数
📖 核心概念¶
1. CUDA线程层次结构¶
Text Only
┌─────────────────────────────────────────────────────────────┐
│ CUDA 线程层次模型 │
├─────────────────────────────────────────────────────────────┤
│ │
│ Grid（网格）← 一次kernel启动产生一个Grid │
│ ├── Block (0,0) Block (1,0) Block (2,0) │
│ ├── Block (0,1) Block (1,1) Block (2,1) │
│ └── Block (0,2) Block (1,2) Block (2,2) │
│ │
│ 每个Block（线程块）包含多个Thread： │
│ Block (1,1) 展开： │
│ ┌────────────────────────────────┐ │
│ │ t(0,0) t(1,0) t(2,0) t(3,0) │ │
│ │ t(0,1) t(1,1) t(2,1) t(3,1) │ ← 最多1024个线程/块 │
│ │ t(0,2) t(1,2) t(2,2) t(3,2) │ │
│ │ t(0,3) t(1,3) t(2,3) t(3,3) │ │
│ └────────────────────────────────┘ │
│ │
│ 层次关系： │
│ Grid → 多个 Block → 每个Block多个 Thread │
│ Grid大小：gridDim.x × gridDim.y × gridDim.z │
│ Block大小：blockDim.x × blockDim.y × blockDim.z │
│ │
└─────────────────────────────────────────────────────────────┘
2. 线程索引计算¶
CUDA
// thread_index.cu - 线程索引计算示例
#include <stdio.h>
#include <cuda_runtime.h>
// 1D Grid + 1D Block
__global__ void kernel_1d(int *output, int n) {
// 全局线程索引
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
output[idx] = idx;
}
}
// 2D Grid + 2D Block（适合矩阵操作）
__global__ void kernel_2d(int *matrix, int width, int height) {
int col = blockIdx.x * blockDim.x + threadIdx.x;
int row = blockIdx.y * blockDim.y + threadIdx.y;
if (col < width && row < height) {
int idx = row * width + col;
matrix[idx] = row * 100 + col; // 存储行列信息
}
}
// 打印每个线程的身份信息
__global__ void identify_threads() {
printf("Grid(%d,%d) Block(%d,%d) Thread(%d,%d) → globalIdx=(%d,%d)\n",
gridDim.x, gridDim.y,
blockIdx.x, blockIdx.y,
threadIdx.x, threadIdx.y,
blockIdx.x * blockDim.x + threadIdx.x,
blockIdx.y * blockDim.y + threadIdx.y);
}
int main() {
// ===== 1D示例 =====
int n = 16;
int *d_output, *h_output;
h_output = (int *)malloc(n * sizeof(int));
cudaMalloc(&d_output, n * sizeof(int));
// 启动配置：4个Block，每Block 4个线程
int blockSize = 4;
int gridSize = (n + blockSize - 1) / blockSize; // 向上取整
printf("1D: gridSize=%d, blockSize=%d\n", gridSize, blockSize);
kernel_1d<<<gridSize, blockSize>>>(d_output, n);
cudaMemcpy(h_output, d_output, n * sizeof(int), cudaMemcpyDeviceToHost);
printf("1D结果: ");
for (int i = 0; i < n; i++) printf("%d ", h_output[i]);
printf("\n\n");
// ===== 2D示例 =====
printf("2D线程身份（2x2 Grid, 2x2 Block）:\n");
dim3 grid2d(2, 2);
dim3 block2d(2, 2);
identify_threads<<<grid2d, block2d>>>();
cudaDeviceSynchronize();
// 清理
free(h_output);
cudaFree(d_output);
return 0;
}
3. 内置变量速查¶
Text Only
┌─────────────────────────────────────────────────────────────┐
│ CUDA 内置变量 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 变量 类型 含义 │
│ ───────────── ────── ────────────────────── │
│ threadIdx.x/y/z uint3 当前线程在Block内的索引 │
│ blockIdx.x/y/z uint3 当前Block在Grid内的索引 │
│ blockDim.x/y/z dim3 Block的维度（每维线程数） │
│ gridDim.x/y/z dim3 Grid的维度（每维Block数） │
│ warpSize int Warp大小（目前固定为32） │
│ │
│ 全局线程ID计算公式（1D）： │
│ globalIdx = blockIdx.x * blockDim.x + threadIdx.x │
│ │
│ 全局线程ID计算公式（2D）： │
│ globalX = blockIdx.x * blockDim.x + threadIdx.x │
│ globalY = blockIdx.y * blockDim.y + threadIdx.y │
│ linearIdx = globalY * (gridDim.x * blockDim.x) + globalX │
│ │
└─────────────────────────────────────────────────────────────┘
4. Warp执行模型¶
Text Only
┌─────────────────────────────────────────────────────────────┐
│ Warp 执行模型 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 一个Block中的线程以Warp（32个线程）为单位执行： │
│ │
│ Block (256 threads) │
│ ├── Warp 0: Thread 0 - 31 ← 32个线程同步执行同一指令 │
│ ├── Warp 1: Thread 32 - 63 │
│ ├── Warp 2: Thread 64 - 95 │
│ ├── ... │
│ └── Warp 7: Thread 224 - 255 │
│ │
│ SIMT（Single Instruction, Multiple Threads）： │
│ 同一Warp中的所有线程在同一时刻执行相同指令 │
│ 但操作各自不同的数据（类似SIMD） │
│ │
│ 关键约束： │
│ • Block大小应为32的倍数（避免不完整Warp浪费） │
│ • 同一Warp内线程执行相同指令路径效率最高 │
│ │
└─────────────────────────────────────────────────────────────┘
5. 分支分歧（Warp Divergence）¶
CUDA
// warp_divergence.cu - 分支分歧对性能的影响
#include <stdio.h>
#include <cuda_runtime.h>
// 有分支分歧的kernel
__global__ void divergent_kernel(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= n) return;
// 同一Warp中的线程走不同分支 → 分歧！
// 偶数线程和奇数线程执行不同代码
if (idx % 2 == 0) {
data[idx] = data[idx] * 2.0f; // 偶数线程执行这里
} else {
data[idx] = data[idx] + 1.0f; // 奇数线程执行这里
}
// 两个分支串行执行，有效吞吐量减半
}
// 无分支分歧的kernel（优化版）
__global__ void no_divergence_kernel(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= n) return;
// 让连续32个线程走同一分支
// 前半Warp全部执行if，后半Warp全部执行else
int warp_id = idx / 32;
if (warp_id % 2 == 0) {
data[idx] = data[idx] * 2.0f; // 整个Warp执行同一分支
} else {
data[idx] = data[idx] + 1.0f;
}
}
int main() {
int n = 1024;
float *d_data;
cudaMalloc(&d_data, n * sizeof(float));
// 初始化数据
float *h_data = (float *)malloc(n * sizeof(float));
for (int i = 0; i < n; i++) h_data[i] = (float)i;
cudaMemcpy(d_data, h_data, n * sizeof(float), cudaMemcpyHostToDevice);
// 启动两个版本对比
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
divergent_kernel<<<gridSize, blockSize>>>(d_data, n);
cudaDeviceSynchronize();
printf("分支分歧演示完成\n");
printf("优化建议：让同一Warp内的线程走相同分支\n");
free(h_data);
cudaFree(d_data);
return 0;
}
6. Kernel启动配置¶
CUDA
// launch_config.cu - 如何选择Grid/Block大小
#include <stdio.h>
#include <cuda_runtime.h>
__global__ void simple_add(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 = 1000000;
size_t size = n * sizeof(float);
// 分配内存
float *h_a = (float *)malloc(size);
float *h_b = (float *)malloc(size);
float *h_c = (float *)malloc(size);
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 初始化
for (int i = 0; i < n; i++) {
h_a[i] = (float)i;
h_b[i] = (float)(i * 2);
}
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 启动配置选择策略
int blockSize = 256; // 常见选择：128, 256, 512
int gridSize = (n + blockSize - 1) / blockSize;
printf("N = %d\n", n);
printf("Block大小: %d线程\n", blockSize);
printf("Grid大小: %d块\n", gridSize);
printf("总线程数: %d（其中%d个多余线程被if保护）\n",
gridSize * blockSize, gridSize * blockSize - n);
simple_add<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
printf("验证: c[0]=%.0f, c[999999]=%.0f\n", h_c[0], h_c[999999]);
// 使用CUDA Occupancy API自动选择blockSize
int minGridSize, optBlockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &optBlockSize,
simple_add, 0, n);
printf("\n推荐配置: blockSize=%d, gridSize=%d\n",
optBlockSize, minGridSize);
// 清理
free(h_a); free(h_b); free(h_c);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}
7. 错误处理¶
CUDA
// error_handling.cu - CUDA错误处理最佳实践
#include <stdio.h>
#include <cuda_runtime.h>
// 错误检查宏
#define CUDA_CHECK(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
fprintf(stderr, "CUDA错误 %s:%d: %s\n", \
__FILE__, __LINE__, cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
// Kernel错误检查
#define KERNEL_CHECK() do { \
cudaError_t err = cudaGetLastError(); \
if (err != cudaSuccess) { \
fprintf(stderr, "Kernel启动错误: %s\n", \
cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
err = cudaDeviceSynchronize(); \
if (err != cudaSuccess) { \
fprintf(stderr, "Kernel执行错误: %s\n", \
cudaGetErrorString(err)); \
exit(EXIT_FAILURE); \
} \
} while(0)
__global__ void demo_kernel(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) data[idx] *= 2.0f;
}
int main() {
float *d_data;
CUDA_CHECK(cudaMalloc(&d_data, 1024 * sizeof(float)));
demo_kernel<<<4, 256>>>(d_data, 1024);
KERNEL_CHECK();
CUDA_CHECK(cudaFree(d_data));
printf("执行成功！\n");
return 0;
}
💡 面试常见问题¶
Q1：Grid、Block、Thread的关系？为什么需要这种层次结构？¶
答：Grid是一次kernel启动的所有线程集合；Block是可以在SM上协作的线程组（共享内存、同步）；Thread是最小执行单元。层次目的：①Block内线程可共享数据和同步；②不同Block可独立调度到不同SM；③便于扩展到不同规模的GPU。
Q2：什么是Warp Divergence？如何避免？¶
答：当同一Warp（32线程）中的线程走不同的if/else分支时，GPU必须串行执行两个分支路径（活跃线程掩码机制），浪费计算资源。避免方法：①让条件基于warp_id而非thread_id；②用数学公式替代分支；③重新组织数据使相邻线程走同一分支。
Q3：如何选择Block大小？¶
答：①必须是32的倍数（避免不完整Warp）；②通常选128/256/512；③考虑寄存器和共享内存限制（影响每SM的活跃Block数）；④可用cudaOccupancyMaxPotentialBlockSizeAPI自动计算；⑤实际性能需Profile确认。
Q4：CUDA的__global__、device、__host__有什么区别？¶
答：__global__是kernel函数，从Host调用在Device执行，返回void；__device__是设备端函数，只能从Device调用在Device执行；__host__是主机端函数（默认）。可以组合__host__ __device__让函数同时为CPU和GPU编译。
Q5：一个Block最多可以有多少线程？受什么限制？¶
答：硬限制1024线程/Block。实际还受限于：①每Block的寄存器用量（总量有限）；②共享内存用量；③这两个资源决定了每个SM能同时运行多少个Block（Occupancy）。可用CUDA Occupancy Calculator分析最优配置。
📝 本章小结¶
Text Only
┌─────────────────────────────────────────────┐
│ 本章核心知识点 │
├─────────────────────────────────────────────┤
│ │
│ 1. Grid→Block→Thread三级层次结构 │
│ 2. 全局线程ID = blockIdx*blockDim+threadIdx│
│ 3. Warp是GPU执行调度的基本单位（32线程） │
│ 4. 分支分歧导致Warp串行执行 │
│ 5. Block大小应为32的倍数 │
│ 6. 始终检查CUDA API调用的错误码 │
│ │
└─────────────────────────────────────────────┘
下一章：03-CUDA编程入门 - 动手编写第一个CUDA程序