03-CUDA编程入门¶
重要性：⭐⭐⭐⭐⭐ 实用度：⭐⭐⭐⭐⭐ 学习时间：3天 必须掌握：是
为什么学这一章？¶
CUDA是NVIDIA推出的并行计算平台和编程模型，是GPU编程的事实标准。掌握CUDA能让你： - 编写高性能的并行程序 - 理解深度学习框架的底层实现 - 优化AI训练和推理性能 - 为学习其他GPU编程技术打下基础
学完这一章，你将能够： - ✅ 编写基本的CUDA程序 - ✅ 理解CUDA内存管理 - ✅ 掌握CUDA线程组织 - ✅ 优化CUDA程序性能
📖 核心概念¶
1. CUDA程序结构¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ CUDA程序基本结构 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ CUDA程序 = CPU代码（Host）+ GPU代码（Device） │
│ │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ CPU代码（Host Code） │ │
│ │ ┌─────────────────────────────────────────────────────┐ │ │
│ │ │ 1. 分配主机内存 │ │ │
│ │ │ float* h_data = new float[N]; │ │ │
│ │ │ │ │ │
│ │ │ 2. 初始化数据 │ │ │
│ │ │ for (int i = 0; i < N; i++) h_data[i] = i; │ │ │
│ │ │ │ │ │
│ │ │ 3. 分配设备内存 │ │ │
│ │ │ float* d_data; │ │ │
│ │ │ cudaMalloc(&d_data, N * sizeof(float)); │ │ │
│ │ │ │ │ │
│ │ │ 4. 拷贝数据到GPU（Host → Device） │ │ │
│ │ │ cudaMemcpy(d_data, h_data, N * sizeof(float), │ │ │
│ │ │ cudaMemcpyHostToDevice); │ │ │
│ │ │ │ │ │
│ │ │ 5. 启动内核（Kernel） │ │ │
│ │ │ myKernel<<<blocks, threads>>>(d_data, N); │ │ │
│ │ │ │ │ │
│ │ │ 6. 等待GPU完成 │ │ │
│ │ │ cudaDeviceSynchronize(); │ │ │
│ │ │ │ │ │
│ │ │ 7. 拷贝结果回CPU（Device → Host） │ │ │
│ │ │ cudaMemcpy(h_data, d_data, N * sizeof(float), │ │ │
│ │ │ cudaMemcpyDeviceToHost); │ │ │
│ │ │ │ │ │
│ │ │ 8. 释放内存 │ │ │
│ │ │ delete[] h_data; │ │ │
│ │ │ cudaFree(d_data); │ │ │
│ │ └─────────────────────────────────────────────────────┘ │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ ↓ │
│ ┌───────────────────────────────────────────────────────────────┐ │
│ │ GPU代码（Device Code） │ │
│ │ ┌─────────────────────────────────────────────────────┐ │ │
│ │ │ __global__ void myKernel(float* data, int n) { │ │ │
│ │ │ // 计算线程ID │ │ │
│ │ │ int i = blockIdx.x * blockDim.x + threadIdx.x; │ │ │
│ │ │ │ │ │
│ │ │ // 检查边界 │ │ │
│ │ │ if (i < n) { │ │ │
│ │ │ // 执行计算 │ │ │
│ │ │ data[i] = data[i] * 2.0f; │ │ │
│ │ │ } │ │ │
│ │ │ } │ │ │
│ │ └─────────────────────────────────────────────────────┘ │ │
│ └───────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────┘
CUDA函数类型修饰符¶
| 修饰符 | 执行位置 | 调用位置 | 说明 |
|---|---|---|---|
__global__ | GPU | CPU | 内核函数，<<< >>>启动 |
__device__ | GPU | GPU | 设备函数，只能被内核调用 |
__host__ | CPU | CPU | 普通函数（默认） |
__host__ __device__ | CPU/GPU | CPU/GPU | 可在两者上编译 |
2. CUDA内存管理¶
内存分配与释放¶
C++
// 主机内存（CPU内存）
float* h_data = new float[N]; // C++方式
cudaMallocHost(&h_data, N * sizeof(float)); // CUDA页锁定内存（更快传输）
delete[] h_data; // 释放C++内存
cudaFreeHost(h_data); // 释放页锁定内存
// 设备内存（GPU显存）
float* d_data;
cudaMalloc(&d_data, N * sizeof(float)); // 分配显存
cudaFree(d_data); // 释放显存
// 统一内存（Unified Memory）- 自动管理
float* u_data;
cudaMallocManaged(&u_data, N * sizeof(float)); // CPU和GPU都能访问
cudaFree(u_data);
内存拷贝¶
C++
// 主机到设备（Host to Device）
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);
// 设备到主机（Device to Host）
cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
// 设备到设备（Device to Device）
cudaMemcpy(d_dst, d_src, size, cudaMemcpyDeviceToDevice);
// 异步拷贝（非阻塞）
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);
内存类型对比¶
Text Only
┌─────────────────────────────────────────────────────────────────────┐
│ CUDA内存类型对比 │
├─────────────────────────────────────────────────────────────────────┤
│ │
│ 1. 页锁定内存（Pinned Memory） │
│ • 使用cudaMallocHost分配 │
│ • 不会被换出到磁盘 │
│ • CPU↔GPU传输速度更快（2倍以上） │
│ • 占用物理内存，不宜分配过多 │
│ │
│ 2. 零拷贝内存（Zero-Copy Memory） │
│ • 使用cudaHostAlloc分配，带cudaHostAllocMapped标志 │
│ • GPU可以直接访问CPU内存 │
│ • 适合小数据量、频繁访问的场景 │
│ • 访问速度慢于显存 │
│ │
│ 3. 统一内存（Unified Memory） │
│ • 使用cudaMallocManaged分配 │
│ • 自动在CPU和GPU间迁移数据 │
│ • 简化编程，但性能可能不如手动管理 │
│ • 适合快速原型开发 │
│ │
│ 性能排序：显存 > 页锁定内存 > 零拷贝内存 > 可分页内存 │
│ │
└─────────────────────────────────────────────────────────────────────┘
3. CUDA线程组织¶
线程索引计算¶
C++
// 一维网格和一维块
__global__ void kernel1D(float* data) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
// idx范围：0 到 (gridDim.x * blockDim.x - 1)
}
// 二维网格和二维块
__global__ void kernel2D(float* data, int width) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
// 适合图像处理
}
// 三维网格和三维块
__global__ void kernel3D(float* data, int width, int height) {
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 * width * height + y * width + x;
// 适合3D模拟
}
线程配置示例¶
C++
// 向量加法（100万元素）
int n = 1000000;
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
// 图像处理（1024x768图像）
dim3 threadsPerBlock(16, 16);
dim3 blocksPerGrid(
(1024 + threadsPerBlock.x - 1) / threadsPerBlock.x,
(768 + threadsPerBlock.y - 1) / threadsPerBlock.y
);
imageProcess<<<blocksPerGrid, threadsPerBlock>>>(d_image, 1024, 768);
4. CUDA优化技术¶
内存访问优化¶
C++
// 不好的内存访问模式（非合并访问）
__global__ void badAccess(float* out, float* in, int n) {
int idx = threadIdx.x;
for (int i = 0; i < n; i += blockDim.x) {
out[i] = in[idx]; // 线程0读in[0], 线程1读in[1]...
// 每个线程访问不同的内存位置，无法合并
}
}
// 好的内存访问模式（合并访问）
__global__ void goodAccess(float* out, float* in, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
out[idx] = in[idx]; // 相邻线程访问相邻内存
// 可以合并成一次内存事务
}
}
共享内存优化¶
C++
// 使用共享内存的矩阵乘法（简化版）
#define TILE_SIZE 16
__global__ void matrixMulShared(float* C, float* A, float* B, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int row = by * TILE_SIZE + ty;
int col = bx * TILE_SIZE + tx;
float sum = 0.0f;
for (int m = 0; m < N / TILE_SIZE; ++m) {
// 协作加载数据到共享内存
As[ty][tx] = A[row * N + m * TILE_SIZE + tx];
Bs[ty][tx] = B[(m * TILE_SIZE + ty) * N + col];
__syncthreads(); // 等待所有线程加载完成
// 计算部分结果
for (int k = 0; k < TILE_SIZE; ++k) {
sum += As[ty][k] * Bs[k][tx];
}
__syncthreads(); // 等待计算完成，防止覆盖数据
}
C[row * N + col] = sum;
}
避免线程发散¶
C++
// 不好的代码（线程发散）
__global__ void divergent(float* out, float* in, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
if (idx % 2 == 0) {
out[idx] = in[idx] * 2; // 偶数线程执行
} else {
out[idx] = in[idx] + 1; // 奇数线程执行
}
// Warp内线程执行不同路径，性能下降
}
}
// 好的代码（避免发散）
__global__ void nonDivergent(float* out, float* in, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float factor = (idx % 2 == 0) ? 2.0f : 0.0f;
float addend = (idx % 2 == 0) ? 0.0f : 1.0f;
out[idx] = in[idx] * factor + addend;
// 所有线程执行相同指令
}
}
🧪 动手实验¶
实验1：矩阵乘法优化¶
目的：学习CUDA优化技术
步骤：
-
创建基础版本
C++// matrix_mul_basic.cu __global__ void matrixMulBasic(float* C, float* A, float* B, int N) { int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if (row < N && col < N) { float sum = 0.0f; for (int k = 0; k < N; k++) { sum += A[row * N + k] * B[k * N + col]; } C[row * N + col] = sum; } } -
创建优化版本（使用共享内存）
C++// matrix_mul_optimized.cu #define TILE_SIZE 32 __global__ void matrixMulOptimized(float* C, float* A, float* B, int N) { __shared__ float As[TILE_SIZE][TILE_SIZE]; __shared__ float Bs[TILE_SIZE][TILE_SIZE]; int bx = blockIdx.x, by = blockIdx.y; int tx = threadIdx.x, ty = threadIdx.y; int row = by * TILE_SIZE + ty; int col = bx * TILE_SIZE + tx; float sum = 0.0f; for (int m = 0; m < (N + TILE_SIZE - 1) / TILE_SIZE; ++m) { if (row < N && m * TILE_SIZE + tx < N) As[ty][tx] = A[row * N + m * TILE_SIZE + tx]; else As[ty][tx] = 0.0f; if (col < N && m * TILE_SIZE + ty < N) Bs[ty][tx] = B[(m * TILE_SIZE + ty) * N + col]; else Bs[ty][tx] = 0.0f; __syncthreads(); for (int k = 0; k < TILE_SIZE; ++k) { sum += As[ty][k] * Bs[k][tx]; } __syncthreads(); } if (row < N && col < N) C[row * N + col] = sum; } -
对比性能
C++int main() { int N = 1024; size_t size = N * N * sizeof(float); // 分配内存、初始化数据... // 测试基础版本 dim3 threads(16, 16); dim3 blocks((N + 15) / 16, (N + 15) / 16); cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start); matrixMulBasic<<<blocks, threads>>>(d_C, d_A, d_B, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float time_basic; cudaEventElapsedTime(&time_basic, start, stop); // 测试优化版本 dim3 threads_opt(32, 32); dim3 blocks_opt((N + 31) / 32, (N + 31) / 32); cudaEventRecord(start); matrixMulOptimized<<<blocks_opt, threads_opt>>>(d_C, d_A, d_B, N); cudaEventRecord(stop); cudaEventSynchronize(stop); float time_optimized; cudaEventElapsedTime(&time_optimized, start, stop); std::cout << "Basic: " << time_basic << " ms" << std::endl; std::cout << "Optimized: " << time_optimized << " ms" << std::endl; std::cout << "Speedup: " << time_basic / time_optimized << "x" << std::endl; return 0; }
实验2：使用CUDA事件计时¶
目的：学习精确的性能测量
步骤：
C++
// timing.cu
#include <cuda_runtime.h>
#include <iostream>
__global__ void workload(float* data, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
float x = data[idx];
for (int i = 0; i < 1000; i++) {
x = sinf(x) + cosf(x);
}
data[idx] = x;
}
}
int main() {
int n = 1000000;
size_t size = n * sizeof(float);
float *d_data;
cudaMalloc(&d_data, size);
// 创建CUDA事件
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// 预热GPU
workload<<<(n + 255) / 256, 256>>>(d_data, n);
cudaDeviceSynchronize();
// 计时
cudaEventRecord(start);
workload<<<(n + 255) / 256, 256>>>(d_data, n);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl;
// 清理
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFree(d_data);
return 0;
}
实验3：使用统一内存¶
目的：简化内存管理
步骤：
C++
// unified_memory.cu
#include <cuda_runtime.h> // 引入头文件
#include <iostream>
__global__ void scale(float* data, int n, float factor) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
data[idx] *= factor;
}
}
int main() {
int n = 1000000;
// 分配统一内存
float* data;
cudaMallocManaged(&data, n * sizeof(float));
// 在CPU上初始化
for (int i = 0; i < n; i++) {
data[i] = i * 0.1f;
}
// 在GPU上执行
scale<<<(n + 255) / 256, 256>>>(data, n, 2.0f);
cudaDeviceSynchronize();
// 在CPU上验证
std::cout << "data[100] = " << data[100] << std::endl;
std::cout << "Expected: " << 100 * 0.1f * 2.0f << std::endl;
// 释放
cudaFree(data);
return 0;
}
💡 核心要点总结¶
CUDA编程流程¶
高级内存优化技术¶
1. Bank Conflict优化¶
共享内存被划分为多个bank，同时访问同一bank的不同地址会导致bank conflict。
C++
// 有Bank Conflict的代码
__global__ void bankConflictExample(float* data) {
__shared__ float shared[256];
int tid = threadIdx.x;
// 32个线程同时访问shared[tid * 2]，产生2-way bank conflict
float val = shared[tid * 2];
}
// 无Bank Conflict的代码
__global__ void noBankConflictExample(float* data) {
__shared__ float shared[256];
int tid = threadIdx.x;
// 32个线程访问连续的bank
float val = shared[tid];
}
解决Bank Conflict的方法： - 使用padding：在共享内存数组中添加额外的列 - 改变访问模式：使用转置或其他访问模式 - 使用广播：当多个线程读取同一地址时，自动广播无conflict
2. 异步数据传输与计算重叠¶
C++
// 使用CUDA Stream实现计算与数据传输重叠
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// 将数据分成两部分，在stream1和stream2中并行处理
for (int i = 0; i < 2; i++) {
int offset = i * halfSize;
cudaMemcpyAsync(d_data + offset, h_data + offset,
halfSize * sizeof(float),
cudaMemcpyHostToDevice,
(i == 0) ? stream1 : stream2);
kernel<<<(halfSize + 255) / 256, 256, 0,
(i == 0) ? stream1 : stream2>>>(d_data + offset);
cudaMemcpyAsync(h_data + offset, d_data + offset,
halfSize * sizeof(float),
cudaMemcpyDeviceToHost,
(i == 0) ? stream1 : stream2);
}
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);
3. 纹理内存和常量内存优化¶
C++
// 常量内存 - 适合只读的小数据（64KB）
__constant__ float constData[256];
cudaMemcpyToSymbol(constData, h_data, 256 * sizeof(float));
// 纹理内存 - 适合具有空间局部性的2D数据
// ⚠️ 注意：以下纹理引用（texture reference）API 在 CUDA 12.0 中已弃用
// 新代码应使用纹理对象（texture object）API，参见 CUDA Programming Guide
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
cudaBindTexture2D(0, texRef, d_data, desc, width, height, pitch);
// 在内核中使用纹理内存
__global__ void textureKernel(float* output, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) {
// 使用纹理缓存，自动处理边界
output[y * width + x] = tex2D(texRef, x, y);
}
}
4. 零拷贝内存（Zero-Copy）¶
C++
// 零拷贝内存 - GPU直接访问CPU内存，适合小数据频繁访问
float* zeroCopyData; // 指针：存储变量的内存地址
cudaHostAlloc(&zeroCopyData, size, cudaHostAllocMapped);
// 获取设备指针
float* d_zeroCopyData;
cudaHostGetDevicePointer(&d_zeroCopyData, zeroCopyData, 0);
// GPU直接访问CPU内存（无需显式拷贝）
kernel<<<blocks, threads>>>(d_zeroCopyData);
优化检查清单¶
- 使用合并内存访问
- 利用共享内存减少全局内存访问
- 避免Bank Conflict
- 避免线程发散
- 选择合适的Block大小（32的倍数）
- 隐藏内存延迟（足够的线程数）
- 使用页锁定内存加速传输
- 考虑使用统一内存简化开发
- 使用CUDA Stream实现计算与传输重叠
- 使用常量内存存储只读小数据
- 使用纹理内存处理2D数据
性能分析工具¶
Bash
# 使用nvprof分析性能（旧版本）
nvprof ./your_program
# 使用Nsight Compute（新版本）
ncu ./your_program
# 查看内存带宽利用率
ncu --metrics dram__bytes.sum.per_second ./your_program
# 查看共享内存bank conflict
ncu --metrics sm__sass_l1_bank_conflict ./your_program
常见错误¶
| 错误 | 原因 | 解决 |
|---|---|---|
| cudaErrorMemoryAllocation | 显存不足 | 减少数据量或分批处理 |
| cudaErrorLaunchFailure | 内核崩溃 | 检查数组越界和非法内存访问 |
| cudaErrorInvalidValue | 参数错误 | 检查线程配置和数据大小 |
| 结果不正确 | 竞争条件 | 使用__syncthreads()同步 |
❓ 常见问题¶
Q1：如何确定最佳的Block大小？
A：一般原则： - 使用Warp大小的倍数（32） - 常见选择：128, 256, 512 - 考虑寄存器使用量（nvcc --ptxas-options=-v查看） - 实验测试不同配置
Q2：为什么我的CUDA程序比CPU还慢？
A：可能原因： - 数据量太小（传输开销 > 计算收益） - 内存访问模式不好（非合并访问） - 线程发散严重 - 没有充分利用并行性
Q3：如何调试CUDA程序？
A：方法： - 使用cuda-gdb（命令行调试器） - 使用Nsight（NVIDIA的IDE插件） - 使用printf在内核中输出 - 检查cudaError_t返回值
Q4：CUDA程序可以移植到AMD GPU吗？
A：不能直接移植。但可以使用： - HIP（AMD的CUDA兼容层） - OpenCL（跨平台） - SYCL（现代C++标准）
📚 扩展阅读¶
- 《CUDA by Example》 - Jason Sanders
- 《Programming Massively Parallel Processors》 - David Kirk
- NVIDIA CUDA最佳实践指南：docs.nvidia.com/cuda/cuda-c-best-practices-guide/
🎯 下一步¶
继续学习GPU并行计算的后续内容，深入理解并行算法设计、内存模型和数据传输优化。