第二章：CUDA编程基础——核函数、线程层次与内存模型

好，咱们直接进入正题。这一章是CUDA编程的基石，说白了就是让你搞明白三件事：核函数怎么写、线程怎么组织、数据怎么放。这三件事搞不清楚，后面做TensorRT优化就是空中楼阁。

2.1 核函数（Kernel）——GPU的入口函数

核函数是什么？简单说，就是跑在GPU上的函数。你写一个函数，前面加个__global__修饰符，然后调用的时候用<<<grid, block>>>这种奇怪的语法——嗯，这就是CUDA编程最标志性的写法。

核函数的基本规则：

  • 必须用 __global__ 修饰
  • 返回类型必须是 void
  • 不能是类的成员函数（静态成员函数可以）
  • 调用时用 <<<Dg, Db, Ns, S>>> 指定执行配置

来看个最简单的例子：

// 核函数：给数组每个元素加1
__global__ void add_one(float *data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        data[idx] += 1.0f;
    }
}

// 调用方式
int n = 1024;
float *d_data;
cudaMalloc(&d_data, n * sizeof(float));
add_one<<<1, 256>>>(d_data, n);  // 1个block，256个线程

这里有个细节我特别想强调：核函数里的变量默认都是寄存器变量，每个线程都有自己的副本。我在项目中遇到过有人把大数组声明在核函数里，结果寄存器爆了，性能直接崩盘。嗯，这个后面讲内存模型时会细说。

2.2 线程层次结构——Grid、Block、Thread

CUDA的线程组织方式，说白了就是三层结构：

  • Thread（线程）：最小的执行单元
  • Block（线程块）：一组线程，可以共享内存、同步
  • Grid（线程网格）：一组Block，构成整个任务

你想想看，这就像军队编制：Grid是师，Block是团，Thread是兵。每个兵有自己的编号，通过threadIdx、blockIdx、blockDim、gridDim这些内置变量来定位。

我的个人习惯：处理一维数组时，我通常用一维的Grid和Block，这样索引计算简单。处理图像时，用二维的Grid和Block，因为图像本身就是二维的，映射起来更自然。

来看个二维的例子：

// 二维核函数：图像灰度化
__global__ void rgb_to_gray(unsigned char *rgb, unsigned char *gray, 
                            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) {
        int idx = (y * width + x) * 3;
        gray[y * width + x] = 0.299f * rgb[idx] + 
                              0.587f * rgb[idx+1] + 
                              0.114f * rgb[idx+2];
    }
}

// 调用方式
dim3 block(16, 16);
dim3 grid((width + 15) / 16, (height + 15) / 16);
rgb_to_gray<<<grid, block>>>(d_rgb, d_gray, width, height);

这里有个坑：Block的线程数不能超过1024（计算能力3.0以上）。我曾经在调试一个模型时，设了dim3 block(32, 32)，结果1024个线程刚好卡在上限，后来换成dim3 block(16, 16)才跑通。所以建议你留点余量，别卡着上限用。

2.3 内存模型简介——数据怎么放是关键

CUDA的内存模型，说白了就是谁能看到什么数据。我刚开始学的时候，觉得这东西很绕，后来发现其实就几个层次：

内存类型 作用域 生命周期 访问速度
寄存器（Register） 单个线程 核函数执行期间 最快
共享内存（Shared Memory） 单个Block Block生命周期 快
全局内存（Global Memory） 所有线程+主机 程序运行期间 慢（约400-800周期）
常量内存（Constant Memory） 所有线程（只读） 程序运行期间 快（有缓存）
纹理内存（Texture Memory） 所有线程（只读） 程序运行期间 适合空间局部性访问

避坑指南：我曾经在优化一个卷积算子时，把所有的中间结果都放到了全局内存里，结果性能还不如CPU。后来改成用共享内存做数据复用，速度直接提升了8倍。记住：能放共享内存的，别放全局内存。

共享内存的使用有个经典模式——分块（Tiling）：

// 矩阵乘法：使用共享内存优化
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    for (int t = 0; t < (N + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        if (row < N && t * TILE_SIZE + threadIdx.x < N)
            As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && t * TILE_SIZE + threadIdx.y < N)
            Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
            
        __syncthreads();  // 同步：确保所有线程都加载完毕
        
        for (int k = 0; k < TILE_SIZE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
            
        __syncthreads();  // 同步：确保所有线程都计算完毕
    }
    
    if (row < N && col < N)
        C[row * N + col] = sum;
}

这段代码里有两个__syncthreads()，我特别提醒一下：同步操作一定要谨慎。如果Block里有些线程提前return了，那同步就会死锁。嗯，这个坑我踩过，调试了整整一天才发现是某个线程提前退出了。

2.4 实战建议——从零开始写第一个核函数

我个人建议你按这个步骤来：

  1. 先写CPU版本：验证算法正确性
  2. 确定线程映射：每个线程处理几个数据？一维还是二维？
  3. 选择Block大小：通常是128或256，要能被32整除（warp大小）
  4. 计算Grid大小：根据数据总量和Block大小计算
  5. 加边界检查：防止越界访问
  6. 测试+调优：用nvprof或nsys分析性能

小技巧：调试核函数时，可以用printf输出信息（计算能力2.0以上支持）。但注意，printf在GPU上很慢，只适合调试用，正式代码要删掉。

好了，这一章的内容就这些。核函数、线程层次、内存模型，这三块是CUDA编程的骨架。下一章我们会深入讲内存访问优化，到时候你会发现——很多时候性能瓶颈不在计算，而在内存。

记住：写CUDA代码，脑子里要时刻想着数据是怎么流动的。数据放对了，性能就来了。

← 上一章 📖 返回目录 下一章 →