跳到主要内容
CUDA编程与算子优化

第1章：CUDA 编程入门

搭建 CUDA 开发环境,理解 Grid/Block/Thread 编程模型和内存模型,编写第一个实用 CUDA Kernel

CUDA 编程模型 内存模型 Kernel

CUDA 是 NVIDIA 提出的并行编程平台,用 C++ 风格的语法描述如何在 GPU 上调度成千上万个线程。本章带你从零搭建开发环境,理解 Grid/Block/Thread 三层线程模型与内存模型,然后从向量加法到矩阵转置写出 4 个完整可运行的 Kernel,把 CUDA 的基本套路过一遍。

📑 目录

1. 环境搭建与最小程序

1.1 安装 CUDA Toolkit

# Ubuntu 22.04
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt update && sudo apt install cuda-toolkit-12-3

# 验证
nvcc --version
nvidia-smi

~/.bashrc 加入:

export PATH=/usr/local/cuda/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH

1.2 第一个 Hello World

// hello.cu
#include <cstdio>

__global__ void hello_kernel() {
    printf("Hello from thread %d in block %d\n",
           threadIdx.x, blockIdx.x);
}

int main() {
    hello_kernel<<<2, 4>>>();   // 2 个 Block,每个 4 个线程,共 8 输出
    cudaDeviceSynchronize();    // 等 GPU 跑完
    return 0;
}

nvcc -arch=sm_80 hello.cu -o hello
./hello

-arch=sm_80 指定为 A100,H100 用 sm_90。架构选错会报”unsupported gpu architecture”。

2. 编程模型:Grid/Block/Thread

2.1 三层结构

Grid    一次 kernel 启动的所有线程
 └── Block  一组线程,共享 Shared Memory,可同步
      └── Thread 最小执行单位

类比:Grid = 整个学校,Block = 一个班级,Thread = 学生。同一个班级的学生可以通过黑板(Shared Memory)交流,跨班级则不行。

2.2 索引计算

// 1D
int idx = blockIdx.x * blockDim.x + threadIdx.x;

// 2D
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

// dim3 启动 2D
dim3 block(16, 16);
dim3 grid((W + 15) / 16, (H + 15) / 16);
my_kernel<<<grid, block>>>(...);

2.3 Block Size 经验法则

场景推荐 Block Size
通用起手256 (= 8 个 Warp)
寄存器紧张128
计算密集 + 共享内存少512-1024

铁律:Block Size 必须是 32 的倍数(与 Warp 对齐)。

2.4 函数修饰符

修饰符调用方执行方用途
__global__CPUGPUKernel,不能有返回值
__device__GPUGPUGPU 上的辅助函数
__host__CPUCPU普通函数(默认)
__host__ __device__双方双方让一份代码 CPU/GPU 都能用

3. 内存模型概览

类型作用域容量速度用途
寄存器线程~256 KB/SM1 cycle私有变量
Shared MemoryBlock~228 KB/SM~30 cycleBlock 内通信
Global Memory(HBM)Grid80 GB~400 cycleKernel 输入输出
ConstantGrid64 KB极快(广播读)不变量
Local线程实际在 HBM慢寄存器溢出

重要:__shared__ 声明的变量住在 Shared Memory,Block 内所有线程共享;局部变量默认住在寄存器,溢出后会落到 Local Memory(其实是 HBM,极慢)。

__global__ void my_kernel(float* data) {
    __shared__ float tile[256];           // Block 共享
    int tid = threadIdx.x;
    tile[tid] = data[blockIdx.x * 256 + tid];
    __syncthreads();                       // Block 内同步
    // ... 用 tile 做计算 ...
}

__syncthreads() 是 Block 内的栅栏——所有线程必须都到这里才能继续。Warp 内可以用更快的 __syncwarp()。

4. 向量加法:你的第一个 Kernel

// vec_add.cu
#include <cstdio>
#include <cstdlib>

__global__ void vec_add(const float* a, const 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() {
    const int N = 1 << 20;     // 1M
    const int bytes = N * sizeof(float);

    // 1. 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);
    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, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. H→D 数据传输
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动 Kernel
    int block = 256;
    int grid = (N + block - 1) / block;
    vec_add<<<grid, block>>>(d_a, d_b, d_c, N);

    // 5. D→H 取结果
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证
    printf("c[1024] = %.0f (expect %.0f)\n", h_c[1024], h_a[1024] + h_b[1024]);

    // 7. 清理
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

🌟 七步套路:分配 host → 分配 device → H→D → kernel → D→H → 验证 → 释放。所有 CUDA 程序的骨架都是这个。

5. 矩阵转置:看到访存模式的力量

5.1 朴素版本(慢)

__global__ void transpose_naive(const float* in, float* out, int W, int H) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    if (x < W && y < H) {
        out[x * H + y] = in[y * W + x];   // 写入是跨 W 步的,非合并访问
    }
}

问题:Warp 内 32 个线程的写入地址相隔 H 个 float,触发 32 次独立 transactions——带宽利用率 1/32。

5.2 Shared Memory 版本(快)

__global__ void transpose_shared(const float* in, float* out, int W, int H) {
    __shared__ float tile[32][33];        // +1 padding 消除 bank conflict
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;

    if (x < W && y < H)
        tile[threadIdx.y][threadIdx.x] = in[y * W + x];   // 合并读

    __syncthreads();

    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    if (x < H && y < W)
        out[y * H + x] = tile[threadIdx.x][threadIdx.y];   // 合并写
}

两个关键技巧:

  1. 借助 Shared Memory 中转:全局内存的读和写都变成合并访问
  2. tile[32][33] 的 +1 padding:消除 bank conflict(下一章详讲)

实测对比(A100,4096×4096):

版本带宽利用率
朴素~40 GB/s (~3%)
Shared Memory~1500 GB/s (~80%)

🌟 核心教训:CUDA 优化的 80% 工作是优化访存模式,不是优化算法。

6. 错误检查与调试

6.1 错误检查宏

#define CUDA_CHECK(call) do {                                    \
    cudaError_t err = call;                                       \
    if (err != cudaSuccess) {                                     \
        fprintf(stderr, "CUDA error %s:%d: %s\n",                 \
                __FILE__, __LINE__, cudaGetErrorString(err));     \
        exit(1);                                                  \
    }                                                             \
} while (0)

CUDA_CHECK(cudaMalloc(&d_a, bytes));

Kernel 错误是异步的——必须 cudaDeviceSynchronize() 后再检查 cudaGetLastError():

my_kernel<<<grid, block>>>(...);
CUDA_CHECK(cudaGetLastError());        // 抓 launch 错误
CUDA_CHECK(cudaDeviceSynchronize());   // 抓 kernel 内部错误

6.2 cuda-memcheck / compute-sanitizer

compute-sanitizer ./my_program          # 检测越界、未初始化、race condition
cuda-gdb ./my_program                   # GPU 上的 gdb

6.3 常见错误

报错原因
invalid configuration argumentBlock size > 1024 或 grid 维度溢出
out of memory显存不够,或 Tensor 没释放
invalid device pointer用了未分配/已释放的 device 指针
unspecified launch failureKernel 内部越界,用 compute-sanitizer 查
unsupported gpu architecture-arch 不匹配实际硬件

✅ 自我检验清单

  • 环境搭建:能从零安装 CUDA Toolkit 并跑通 Hello World
  • 三层结构:能向小白解释 Grid / Block / Thread 的关系,以及为什么 Block Size 要 32 的倍数
  • 索引计算:能正确写出 1D/2D 全局线程索引
  • 内存层次:能解释寄存器、Shared、Global、Constant 各自的特点和用途
  • 向量加法:能不看资料默写完整的 vec_add 程序
  • 矩阵转置:能解释朴素版为什么慢、Shared Memory 版为什么快、+1 padding 解决了什么
  • 错误检查:能写出 CUDA_CHECK 宏并解释为什么 Kernel 错误是异步的
  • compute-sanitizer:遇到 unspecified launch failure,能用工具定位到具体哪行越界

📚 参考资料