CUDA系统学习教程
CUDA 系统学习教程
课程大纲
| 课次 | 主题 | 重点内容 |
|---|---|---|
| 1 | CUDA 基础概念 | GPU 架构、异构计算模型 |
| 2 | 线程层级结构 | Grid、Block、Thread |
| 3 | 内核函数 | __global__、__device__、启动语法 |
| 4 | 内存管理 | cudaMalloc、cudaMemcpy、cudaFree |
| 5 | 线程索引计算 | blockIdx、threadIdx、多维索引 |
| 6 | 并行计算模式 | 向量加法、矩阵乘法 |
| 7 | 同步与共享内存 | __shared__、__syncthreads() |
| 8 | 性能优化 | 内存合并、避免分支分歧 |
第一课：CUDA 基础概念
知识点
什么是 CUDA？
CUDA = Compute Unified Device Architecture（统一计算设备架构）
- NVIDIA 开发的并行计算平台和编程模型
- 让程序员可以使用 C/C++ 直接编程 GPU
- 实现 CPU（主机）和 GPU（设备）的协同工作
异构计算模型
┌─────────────────────────────────────────────────┐
│ 程序执行流程 │
├─────────────────────────────────────────────────┤
│ │
│ CPU（主机） GPU（设备） │
│ ├─ 串行代码 ├─ 并行代码 │
│ ├─ 逻辑控制 ├─ 大量计算 │
│ └─ 内存管理 └─ 数据并行 │
│ │
│ CPU 准备数据 ──→ 传输到 GPU ──→ GPU 计算 │
│ CPU 接收结果 ←── 传回 CPU ←── GPU 完成 │
│ │
└─────────────────────────────────────────────────┘
CPU vs GPU 对比
| 特性 | CPU | GPU |
|---|---|---|
| 核心数 | 少（4-64） | 多（数千） |
| 时钟频率 | 高（3-5 GHz） | 低（1-2 GHz） |
| 缓存 | 大 | 小 |
| 适用场景 | 串行、逻辑控制 | 并行、数值计算 |
| 内存 | 主机内存（Host Memory） | 设备内存（Device Memory） |
CUDA 程序基本结构
int main() {
// 1. 分配主机内存
float *h_data = (float*)malloc(size);
// 2. 分配设备内存
float *d_data;
cudaMalloc(&d_data, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 4. 启动内核（并行计算）
myKernel<<<grid, block>>>(d_data);
// 5. 拷贝结果回主机
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
// 6. 释放内存
cudaFree(d_data);
free(h_data);
return 0;
}
练习题 1
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | CUDA 的全称是什么？ | Compute Unified Device Architecture（统一计算设备架构） |
| 2 | CPU 和 GPU 分别适合什么类型的任务？ | CPU：串行、逻辑控制、复杂决策 GPU：并行、数值计算、大规模数据处理 |
| 3 | CUDA 程序的基本执行流程？ | 1. 分配内存（主机+设备） 2. 数据传输到设备 3. 启动内核并行计算 4. 结果传回主机 5. 释放内存 |
第二课：线程层级结构
知识点
三级线程层级
CUDA 使用三级线程层级来组织并行执行：
┌─────────────────────────────────────────────────┐
│ Grid（网格） │
│ 一个内核启动产生一个 Grid │
│ │
│ ┌─────────────┐ ┌─────────────┐ ┌─────────┐ │
│ │ Block 0 │ │ Block 1 │ │ Block N │ │
│ │ 线程块 │ │ 线程块 │ │ 线程块 │ │
│ │ │ │ │ │ │ │
│ │ T0 T1 T2... │ │ T0 T1 T2... │ │ T0 T1.. │ │
│ │ 线程 │ │ 线程 │ │ 线程 │ │
│ └─────────────┘ └─────────────┘ └─────────┘ │
│ │
│ Grid = 所有 Block 的集合 │
│ Block = 一组 Thread 的集合 │
│ Thread = 最小执行单元 │
└─────────────────────────────────────────────────┘
层级说明
| 层级 | 说明 | 数量限制 |
|---|---|---|
| Grid | 网格，一个内核启动的所有线程 | 最多 2^31-1 个 Block |
| Block | 线程块，一组可协作的线程 | 最多 1024 个 Thread |
| Thread | 线程，最小执行单元 | 执行一个内核函数实例 |
Block 和 Grid 的维度
可以是 一维、二维或三维：
// 一维
dim3 grid(10); // 10 个 Block
dim3 block(256); // 每个 Block 256 个 Thread
// 二维
dim3 grid(10, 10); // 10x10 = 100 个 Block
dim3 block(16, 16); // 每个 Block 16x16 = 256 个 Thread
// 三维
dim3 grid(10, 10, 10); // 10x10x10 = 1000 个 Block
dim3 block(8, 8, 8); // 每个 Block 8x8x8 = 512 个 Thread
为什么使用多维？
- 一维：处理数组、向量
- 二维：处理图像、矩阵
- 三维：处理体积数据、3D 模型
练习题 2
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | CUDA 的三级线程层级是什么？ | Grid（网格）→ Block（线程块）→ Thread（线程） |
| 2 | 一个 Block 最多能有多少个 Thread？ | 1024 个 |
| 3 | 为什么需要多维的 Block 和 Grid？ | 不同维度的数据结构：一维处理数组，二维处理图像/矩阵，三维处理体积数据 |
第三课：内核函数
知识点
函数修饰符
CUDA 使用特殊修饰符区分不同类型的函数：
| 修饰符 | 执行位置 | 调用位置 | 说明 |
|---|---|---|---|
__global__ |
GPU | CPU 或 GPU | 内核函数，启动并行执行 |
__device__ |
GPU | GPU | 设备函数，只能被 GPU 调用 |
__host__ |
CPU | CPU | 主机函数，普通 C/C++ 函数 |
内核函数定义
// __global__ 内核函数
__global__ void myKernel(int *data, int n) {
int idx = threadIdx.x; // 获取线程索引
if (idx < n) {
data[idx] = data[idx] * 2; // 每个线程处理一个元素
}
}
内核启动语法
// 启动内核
myKernel<<<gridSize, blockSize>>>(参数列表);
// gridSize: Block 数量（可以是 int 或 dim3）
// blockSize: 每个 Block 的 Thread 数量（可以是 int 或 dim3）
示例
// 一维启动
int gridSize = 10; // 10 个 Block
int blockSize = 256; // 每个 Block 256 个 Thread
myKernel<<<gridSize, blockSize>>>(d_data, n);
// 二维启动
dim3 grid(10, 10); // 10x10 个 Block
dim3 block(16, 16); // 每个 Block 16x16 个 Thread
myKernel<<<grid, block>>>(d_matrix, width, height);
内核函数限制
- 不能有返回值（必须是 void）
- 不能使用递归
- 不能使用静态变量
- 不能使用可变参数
练习题 3
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | __global__ 修饰符的作用是什么？ |
定义内核函数，在 GPU 上执行，可从 CPU 或 GPU 调用 |
| 2 | 写出启动内核 addKernel 的语法，使用 5 个 Block，每个 Block 128 个 Thread |
addKernel<<<5, 128>>>(参数); |
| 3 | 内核函数有哪些限制？ | 不能有返回值、不能递归、不能使用静态变量、不能使用可变参数 |
第四课：内存管理
知识点
CUDA 内存类型
| 内存类型 | 位置 | 作用 |
|---|---|---|
| 主机内存（Host Memory） | CPU | 存储输入数据和接收结果 |
| 设备内存（Device Memory） | GPU | 存储 GPU 计算数据 |
| 共享内存（Shared Memory） | GPU Block 内 | Block 内线程共享，高速 |
| 常量内存（Constant Memory） | GPU | 只读，缓存优化 |
| 全局内存（Global Memory） | GPU | 所有线程可访问 |
基本内存操作函数
// 1. 分配设备内存
cudaMalloc(void **ptr, size_t size);
// 2. 拷贝数据
cudaMemcpy(void *dst, void *src, size_t size, cudaMemcpyKind kind);
// 3. 释放设备内存
cudaFree(void *ptr);
cudaMemcpyKind 类型
| 类型 | 方向 |
|---|---|
| cudaMemcpyHostToDevice | 主机 → 设备 |
| cudaMemcpyDeviceToHost | 设备 → 主机 |
| cudaMemcpyDeviceToDevice | 设备 → 设备 |
完整示例
int main() {
int n = 1000;
size_t size = n * sizeof(float);
// 1. 分配主机内存
float *h_a = (float*)malloc(size);
float *h_b = (float*)malloc(size);
float *h_c = (float*)malloc(size);
// 初始化数据
for (int i = 0; i < n; i++) {
h_a[i] = i;
h_b[i] = i * 2;
}
// 2. 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 4. 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 5. 拷贝结果回主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 6. 释放内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(h_a);
free(h_b);
free(h_c);
return 0;
}
错误检查
// 检查 CUDA 错误
cudaError_t err = cudaMalloc(&d_data, size);
if (err != cudaSuccess) {
printf("CUDA 错误: %s\n", cudaGetErrorString(err));
}
练习题 4
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | 写出分配 1000 个 float 的设备内存的代码 | float *d_data; cudaMalloc(&d_data, 1000 * sizeof(float)); |
| 2 | cudaMemcpy 的四个参数分别是什么？ | 目标地址、源地址、数据大小、传输方向 |
| 3 | 为什么需要 cudaFree？ | 释放设备内存，避免内存泄漏 |
第五课：线程索引计算
知识点
内置变量
CUDA 提供内置变量来获取线程索引：
| 变量 | 说明 | 类型 |
|---|---|---|
| blockIdx | 当前 Block 在 Grid 中的索引 | dim3 |
| threadIdx | 当前 Thread 在 Block 中的索引 | dim3 |
| blockDim | Block 的维度（Thread 数量） | dim3 |
| gridDim | Grid 的维度（Block 数量） | dim3 |
一维索引计算
// 全局线程索引（一维）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 示例：
// blockIdx.x = 2, blockDim.x = 256, threadIdx.x = 10
// idx = 2 * 256 + 10 = 522
二维索引计算
// 全局线程索引（二维）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
// 转换为一维索引（处理矩阵）
int idx = y * width + x;
三维索引计算
// 全局线程索引（三维）
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
// 转换为一维索引
int idx = z * height * width + y * width + x;
边界检查
__global__ void myKernel(float *data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// 必须检查边界！
if (idx < n) {
data[idx] = data[idx] * 2;
}
}
为什么需要边界检查？
- 数据大小可能不是 blockSize 的整数倍
- 多余的线程不应该访问无效内存
计算 Grid 大小
// 确保 Grid 大小足够覆盖所有数据
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize; // 向上取整
// 示例：
// n = 1000, blockSize = 256
// gridSize = (1000 + 255) / 256 = 4
// 总线程数 = 4 * 256 = 1024 > 1000 ✓
练习题 5
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | 写出一维全局线程索引的计算公式 | int idx = blockIdx.x * blockDim.x + threadIdx.x; |
| 2 | blockIdx.x=3, blockDim.x=128, threadIdx.x=50，全局索引是多少？ | idx = 3 * 128 + 50 = 434 |
| 3 | 为什么需要边界检查？ | 数据大小可能不是 blockSize 的整数倍，多余线程不应访问无效内存 |
第六课：并行计算模式
知识点
向量加法
// 内核函数
__global__ void vectorAdd(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 = 1000;
size_t size = n * sizeof(float);
// 分配和初始化主机内存...
// 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 拷贝数据
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 拷贝结果
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 释放内存...
}
矩阵乘法
// 内核函数（简单版本）
__global__ void matrixMul(float *A, float *B, float *C, int width) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < width && col < width) {
float sum = 0.0f;
for (int k = 0; k < width; k++) {
sum += A[row * width + k] * B[k * width + col];
}
C[row * width + col] = sum;
}
}
// 启动内核
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (width + 15) / 16);
matrixMul<<<grid, block>>>(d_A, d_B, d_C, width);
并行计算模式总结
| 模式 | 特点 | 应用场景 |
|---|---|---|
| 元素级并行 | 每个线程处理一个元素 | 向量加法、标量乘法 |
| 行/列并行 | 每个线程处理一行/列 | 矩阵操作 |
| 归约 | 多线程协作合并结果 | 求和、最大值、最小值 |
| 扫描 | 计算前缀和 | 排序、压缩 |
练习题 6
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | 向量加法中，每个线程做什么？ | 处理一个元素：c[idx] = a[idx] + b[idx] |
| 2 | 矩阵乘法中，row 和 col 如何计算？ | row = blockIdx.y * blockDim.y + threadIdx.y col = blockIdx.x * blockDim.x + threadIdx.x |
| 3 | 为什么矩阵乘法使用二维 Block？ | 矩阵是二维结构，二维 Block 更直观地映射到矩阵元素 |
第七课：同步与共享内存
知识点
共享内存
共享内存是 Block 内所有线程共享的高速内存：
__global__ void myKernel(float *data) {
// 声明共享内存
__shared__ float sharedData[256];
int idx = threadIdx.x;
// 每个线程加载一个元素到共享内存
sharedData[idx] = data[idx];
// 同步，确保所有线程都完成加载
__syncthreads();
// 使用共享内存计算
data[idx] = sharedData[idx] * 2;
}
共享内存特点
| 特性 | 说明 |
|---|---|
| 位置 | GPU 片上，位于 Block 内 |
| 速度 | 比全局内存快约 100 倍 |
| 大小 | 每个 Block 最多 48KB（可配置） |
| 生命周期 | Block 执行期间 |
| 可见性 | 仅 Block 内线程可见 |
同步函数
__syncthreads(); // Block 内所有线程同步
作用：
- 确保所有线程都执行到同一位置
- 用于共享内存读写同步
注意：
- 只能在 Block 内同步
- 不能在条件分支中调用（可能导致死锁）
共享内存应用：矩阵乘法优化
__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
__shared__ float As[16][16];
__shared__ float Bs[16][16];
int row = blockIdx.y * 16 + threadIdx.y;
int col = blockIdx.x * 16 + threadIdx.x;
float sum = 0.0f;
// 分块计算
for (int k = 0; k < width; k += 16) {
// 加载到共享内存
As[threadIdx.y][threadIdx.x] = A[row * width + k + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * width + col];
__syncthreads();
// 计算部分结果
for (int i = 0; i < 16; i++) {
sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
}
__syncthreads();
}
C[row * width + col] = sum;
}
练习题 7
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | 共享内存用什么关键字声明？ | shared |
| 2 | 共享内存比全局内存快多少倍？ | 约 100 倍 |
| 3 | __syncthreads() 的作用是什么？ | Block 内所有线程同步，确保都执行到同一位置 |
| 4 | 为什么 __syncthreads() 不能在条件分支中调用？ | 可能导致部分线程跳过同步，造成死锁 |
第八课：性能优化
知识点
1. 内存合并（Memory Coalescing）
概念：相邻线程访问相邻内存地址时，GPU 可以合并为一次内存访问。
// 好的访问模式（合并）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
data[idx] = value; // 线程 0 访问地址 0，线程 1 访问地址 1...
// 坏的访问模式（不合并）
int idx = threadIdx.x * stride; // 线程 0 访问地址 0，线程 1 访问地址 stride...
data[idx] = value;
2. 避免分支分歧（Branch Divergence）
概念：同一个 Warp（32 个线程）内的线程执行不同分支时，会串行执行。
// 坏的分支（分歧）
if (threadIdx.x % 2 == 0) {
// 偶数线程执行
} else {
// 奇数线程执行
}
// 好的分支（无分歧）
if (threadIdx.x < 16) {
// 前 16 个线程执行
} else {
// 后 16 个线程执行
}
3. 使用共享内存减少全局内存访问
// 不使用共享内存（多次访问全局内存）
for (int i = 0; i < N; i++) {
sum += data[idx + i]; // 每次都访问全局内存
}
// 使用共享内存（一次加载，多次使用）
__shared__ float sharedData[BLOCK_SIZE];
sharedData[threadIdx.x] = data[idx];
__syncthreads();
for (int i = 0; i < N; i++) {
sum += sharedData[i]; // 访问共享内存
}
4. 合适的 Block 大小
- 通常是 128、256 或 512
- 需要考虑：
- 寄存器使用量
- 共享内存使用量
- Warp 数量（建议至少 2 个 Warp = 64 个线程）
5. 使用 CUDA 事件计时
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
myKernel<<<grid, block>>>(...);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("执行时间: %.2f ms\n", milliseconds);
cudaEventDestroy(start);
cudaEventDestroy(stop);
性能优化总结
| 优化方法 | 说明 | 效果 |
|---|---|---|
| 内存合并 | 相邻线程访问相邻地址 | 提高内存带宽利用率 |
| 避免分支分歧 | Warp 内线程执行相同分支 | 减少串行执行 |
| 使用共享内存 | 减少全局内存访问 | 大幅提升速度 |
| 合适 Block 大小 | 128/256/512 | 平衡资源使用 |
| 减少数据传输 | 最小化 Host-Device 传输 | 减少传输延迟 |
练习题 8
| 题号 | 问题 | 正确答案 |
|---|---|---|
| 1 | 什么是内存合并？ | 相邻线程访问相邻内存地址时，GPU 合并为一次访问 |
| 2 | Warp 的大小是多少？ | 32 个线程 |
| 3 | 为什么分支分歧会影响性能？ | Warp 内线程执行不同分支时会串行执行，降低并行效率 |
| 4 | 共享内存为什么能提高性能？ | 比全局内存快约 100 倍，减少全局内存访问次数 |
综合练习：完整 CUDA 程序
向量加法完整代码
#include <stdio.h>
// 内核函数
__global__ void vectorAdd(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 = 10000;
size_t size = n * sizeof(float);
// 1. 分配主机内存
float *h_a = (float*)malloc(size);
float *h_b = (float*)malloc(size);
float *h_c = (float*)malloc(size);
// 初始化数据
for (int i = 0; i < n; i++) {
h_a[i] = (float)i;
h_b[i] = (float)(i * 2);
}
// 2. 分配设备内存
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// 3. 拷贝数据到设备
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
// 4. 启动内核
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);
// 5. 拷贝结果回主机
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
// 6. 验证结果
bool success = true;
for (int i = 0; i < n; i++) {
if (h_c[i] != h_a[i] + h_b[i]) {
success = false;
break;
}
}
printf("验证结果: %s\n", success ? "成功" : "失败");
// 7. 释放内存
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
free(h_a);
free(h_b);
free(h_c);
return 0;
}
学习检查清单
基础概念 ✅
线程组织 ✅
内核函数 ✅
内存管理 ✅
线程索引 ✅
并行计算 ✅
同步与共享内存 ✅
性能优化 ✅
下一步学习建议
- 动手实践：运行教程中的示例代码
- 参考 CUDA_Freshman 目录：查看更多示例
- 阅读 NVIDIA 文档：CUDA C Programming Guide
- 使用 CUDA Profiler：nsight systems、nsight compute
学习日期：2026-05-07
浙公网安备 33010602011771号