CUDA 使用示例与性能剖析手段

CUDA 使用示例与性能剖析手段

示例一：向量加法，展示最基本的 kernel 写法与执行配置，注意边界检查与页锁定内存。

#include <cuda_runtime.h>

__global__ void vec_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() {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);

    float *ha, *hb, *hc;
    cudaMallocHost(&ha, bytes);   // 页锁定内存，配合异步拷贝
    cudaMallocHost(&hb, bytes);
    cudaMallocHost(&hc, bytes);

    float *da, *db, *dc;
    cudaMalloc(&da, bytes);
    cudaMalloc(&db, bytes);
    cudaMalloc(&dc, bytes);

    cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    vec_add<<<blocks, threads>>>(da, db, dc, n);

    cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();

    cudaFree(da); cudaFree(db); cudaFree(dc);
    cudaFreeHost(ha); cudaFreeHost(hb); cudaFreeHost(hc);
    return 0;
}

示例二：用 shared memory 分块做矩阵乘法，把 global memory 的重复读取降到 O(1/TILE)，同时避免 bank conflict。

#define TILE 16

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

    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;
    int col = blockIdx.x * TILE + tx;
    float acc = 0.0f;

    for (int t = 0; t < (k + TILE - 1) / TILE; ++t) {
        As[ty][tx] = (row < m && t * TILE + tx < k)
                   ? A[row * k + t * TILE + tx] : 0.0f;
        Bs[ty][tx] = (t * TILE + ty < k && col < n)
                   ? B[(t * TILE + ty) * n + col] : 0.0f;
        __syncthreads();
        for (int i = 0; i < TILE; ++i) {
            acc += As[ty][i] * Bs[i][tx];
        }
        __syncthreads();
    }
    if (row < m && col < n) {
        C[row * n + col] = acc;
    }
}

编译与剖析：

# 编译，-lineinfo 让剖析工具能把指标对应回源码行
nvcc -O3 -arch=native -lineinfo matmul.cu -o matmul

# 系统级时间线：看 kernel 之间是否有空隙、拷贝与计算是否重叠
nsys profile --stats=true ./matmul

# 单 kernel 深挖：访存吞吐、占用率、stall 原因、指令混合
ncu --set full --launch-count 3 ./matmul
ncu --metrics gpu__time_duration.sum,dram__throughput.avg.pct_of_peak_sustained_elapsed ./matmul

看指标时抓几个关键量：dram__throughput 接近峰值说明卡在带宽上，这时提占用率没有收益，应该改数据复用；achieved occupancy 低且 stall 主要由长延迟访存构成，就该调 block 大小或降寄存器用量；smsp__inst_executed 里分支类指令占比高，则要处理发散。另一个必备工具是 compute-sanitizer，用 compute-sanitizer ./matmul 抓越界访问、未初始化读取和竞态。

适合用 CUDA 手写的场景：算子结构规整、数据并行度高、访存模式可预测，比如归一化、激活、量化与反量化、注意力里的重排、以及各种融合算子。不适合的场景：分支复杂且数据相关的算法（排序、图遍历）、依赖大量串行逻辑的任务、以及矩阵乘这类已经由 cuBLAS / CUTLASS 高度优化过的运算——自己写通常更慢。