cuda

CUDA

目的

指导 Agent 进行 NVIDIA CUDA C/C++ 开发：内核启动配置、从寄存器到全局内存的存储层次结构、使用流进行异步执行、nvcc 编译标志、Thrust 库的使用，以及诊断常见性能问题，例如 warp 分歧和非合并内存访问。

适用场景

• 为矩阵乘法、归约或模板运算编写或优化 CUDA 内核
• 选择 block/grid 维度并估算占用率
• 调试因内存访问模式或占用率低而运行缓慢的内核
• 使用异步 cudaMemcpyAsync 设置多流流水线
• 使用 nvcc 编译并选择架构标志（-gencode）
• 使用 Thrust 执行并行原语，而不是手写内核

工作流

1. 最小内核与启动

CUDA 将工作组织为线程组（block），多个 block 组成一个 grid。

Thread hierarchy
├── grid (1D/2D/3D)
│   └── block (1D/2D/3D, max 1024 threads)
│       └── thread (threadIdx, blockIdx, blockDim, gridDim)

// vector_add.cu
#include <cuda_runtime.h>
#include <stdio.h>

__global__ void vector_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 main(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *h_a, *h_b, *h_c, *d_a, *d_b, *d_c;

    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    // ... host init and cudaMemcpy H2D ...

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vector_add<<<blocks, threads>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

2. 存储层次结构

存储类型	作用范围	延迟	典型用途
寄存器	每个线程	~1 个周期	局部标量、循环索引
共享内存（__shared__）	每个 block	~5 个周期	分块数据、halo 交换
全局内存	所有线程	~400+ 个周期	大型数组、合并访问
常量内存（__constant__）	只读、缓存	快速广播	内核参数、查找表
纹理内存	缓存的二维访问	已缓存	图像采样、不规则读取

共享内存示例（矩阵分块）：

#define TILE 16

__global__ void matmul_tiled(const float *A, const float *B, float *C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
        As[threadIdx.y][threadIdx.x] = (row < N && t * TILE + threadIdx.x < N)
            ? A[row * N + t * TILE + threadIdx.x] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (col < N && t * TILE + threadIdx.y < N)
            ? B[(t * TILE + threadIdx.y) * N + col] : 0.0f;
        __syncthreads();

        for (int k = 0; k < TILE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < N && col < N)
        C[row * N + col] = sum;
}

3. 流与异步拷贝

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, h_b, bytes, cudaMemcpyHostToDevice, stream2);
vector_add<<<blocks, threads, 0, stream1>>>(d_a, d_b, d_c, n);
cudaMemcpyAsync(h_c, d_c, bytes, cudaMemcpyDeviceToHost, stream1);
cudaStreamSynchronize(stream1);

锁页主机内存（cudaMallocHost）可实现真正的异步 DMA 与内核执行重叠。

4. nvcc 编译

# Single architecture (local GPU)
nvcc -O3 -arch=sm_80 -o prog vector_add.cu

# Fat binary for multiple GPUs
nvcc -O3 \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -o prog vector_add.cu

# Debug symbols for cuda-gdb
nvcc -G -g -O0 -arch=sm_80 -o prog_debug vector_add.cu

# Show PTX/SASS
nvcc -arch=sm_80 -ptx vector_add.cu
nvcc -arch=sm_80 -cubin vector_add.cu
cuobjdump -sass prog

常用标志：

标志	效果
-O3	激进优化
-G	调试时禁用优化
-lineinfo	在性能分析中关联源代码行
-Xcompiler -fopenmp	在 CUDA 中使用主机端 OpenMP
--use_fast_math	使用更快、精度较低的数学内建函数
-maxrregcount=N	限制寄存器数量以提高占用率

5. 占用率估算

# CUDA Occupancy Calculator (spreadsheet) or programmatic:
./occupancy_tool --kernel vector_add --block-size 256 --regs 16 --smem 0

决策树：

Kernel slow?
├── Low occupancy (< 25%) → reduce registers, shared mem, or block size
├── Memory-bound → check coalescing, use shared memory tiling
└── Compute-bound → increase arithmetic intensity, use tensor cores

使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor API，或 Nsight Compute 的 sm__warps_active.avg.pct_of_peak_sustained_active 指标。

6. Thrust 基础

#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/reduce.h>

thrust::device_vector<int> d_vec(1000000);
thrust::sort(d_vec.begin(), d_vec.end());
int sum = thrust::reduce(d_vec.begin(), d_vec.end());

Thrust 会在内部处理临时存储和内核启动。排序、扫描和归约优先使用 Thrust；对于特定领域的融合操作，则编写自定义内核。

7. 常见问题

Warp 分歧：一个 warp 中的线程（32 个）以 SIMT 锁步方式执行。根据 threadIdx 进行分支会导致串行执行。

// Bad: divergent branch
if (threadIdx.x % 2 == 0) { heavy_a(); } else { heavy_b(); }

// Better: separate kernels or predication

非合并访问：相邻线程应访问相邻地址。

// Bad: strided access
float val = data[threadIdx.x * stride];

// Good: coalesced
float val = data[blockIdx.x * blockDim.x + threadIdx.x];

常见问题

现象	原因	修复方法
cudaErrorIllegalAddress（700）	设备指针越界	运行 compute-sanitizer --tool memcheck；添加边界检查
cudaErrorLaunchTimeout（702）	内核运行时间超过看门狗限制	减少工作量；拆分内核；在开发用 GPU 上禁用 TDR
占用率低警告	寄存器或共享内存使用过多	减少 __shared__ 大小；使用 -maxrregcount；缩小 block
CPU 与 GPU 计算结果不一致	共享内存或全局内存上的竞争	添加 __syncthreads()；在归约中使用原子操作
no kernel image available	-arch=sm_XX 不正确	与 GPU 计算能力匹配：nvidia-smi --query-gpu=compute_cap
H2D/D2H 拷贝缓慢	主机内存可分页	使用 cudaMallocHost 分配锁页缓冲区

相关技能

• skills/gpu/cuda-profiling — 使用 Nsight Systems/Compute 诊断内核性能
• skills/gpu/cuda-debugging — 使用 cuda-gdb 和 compute-sanitizer 检查正确性
• skills/gpu/gpu-memory-model — SIMT 执行、合并规则、bank 冲突
• skills/gpu/hip-rocm — 将 CUDA 内核移植到 AMD HIP
• skills/gpu/triton-lang — 用于自定义内核的 Python DSL 替代方案
• skills/hpc/openmp — 与 CUDA 并用的主机端并行处理