3. CUDA编程核心概念

好,咱们今天聊点硬核的。CUDA编程的核心,说白了就三件事:线程怎么组织、数据怎么放、函数怎么写。这三件事搞明白了,你写出来的Kernel性能就不会太差。

我记得刚接触CUDA那会儿,看了半天文档,感觉每个概念都懂,一上手写代码就懵。后来踩了无数坑才明白——线程层次结构和内存层次结构,是CUDA性能的基石。

3.1 线程层次结构:Grid / Block / Thread

先说说线程怎么组织。CUDA的线程模型是三级结构:

  • Thread(线程):最基础的执行单元
  • Block(线程块):一组线程,可以协作
  • Grid(网格):一组Block,构成整个任务

你想想看,这就像军队编制:Thread是士兵,Block是班,Grid是连。每个士兵独立作战,但同一个班的士兵可以互相配合。

关键点:同一个Block内的线程可以通过Shared Memory通信,不同Block的线程不能直接通信。

我在项目中遇到过一个问题:有个同事把数据分成了1024个Block,每个Block只放1个线程。结果呢?性能惨不忍睹。为什么?因为GPU的调度单位是Warp(32个线程),一个Block里只有1个线程,意味着31个CUDA核心在空转。

来看个实际例子:

// 定义线程层次结构
dim3 gridDim(4, 4);      // 4x4 = 16个Block
dim3 blockDim(16, 16);   // 16x16 = 256个线程/Block

// 启动Kernel
myKernel<<<gridDim, blockDim>>>(d_data, N);

这里总共有16 × 256 = 4096个线程。每个线程通过内置变量知道自己是谁:

  • threadIdx.x:线程在Block内的索引
  • blockIdx.x:Block在Grid内的索引
  • blockDim.x:Block的维度大小

计算全局索引的公式很简单:

int tid = blockIdx.x * blockDim.x + threadIdx.x;

我的习惯:Block大小通常设为128或256的倍数。为什么?因为Warp是32线程,128正好是4个Warp,调度效率最高。

3.2 内存层次结构:Global / Shared / Register

内存这块,是CUDA性能优化的重头戏。我见过太多人把数据全扔在Global Memory里,然后抱怨GPU慢。

CUDA的内存层次,从慢到快依次是:

内存类型 位置 访问速度 生命周期 作用域
Global Memory 显存 ~400-800周期 程序运行期间 所有线程
Shared Memory 芯片内 ~5-10周期 Block生命周期 同一Block
Register 寄存器 ~1周期 线程生命周期 单个线程

说白了,Global Memory就是仓库,Shared Memory是工作台,Register是你手里的工具。你总不能每次拧螺丝都跑一趟仓库拿扳手吧?

我曾经踩过的坑:有一次写矩阵乘法,直接把两个大矩阵都放在Global Memory里,每次计算都去读。结果性能只有理论峰值的5%。后来改成用Shared Memory做分块,性能直接翻了10倍。

Shared Memory的使用方式:

__global__ void kernel() {
    // 声明Shared Memory
    __shared__ float s_data[256];
    
    // 从Global加载到Shared
    s_data[threadIdx.x] = global_data[blockIdx.x * blockDim.x + threadIdx.x];
    
    // 同步,确保所有线程加载完成
    __syncthreads();
    
    // 从Shared Memory读取,速度飞快
    float val = s_data[threadIdx.x];
}

这里有个细节:__syncthreads()是Block内的屏障。所有线程必须执行到这里,才能继续往下走。少了它,数据可能还没加载完就被读了,结果就是随机错误。

3.3 Kernel函数编写

Kernel函数是跑在GPU上的代码。写法上跟CPU函数差不多,但有几个关键区别:

  • 用__global__修饰符
  • 返回类型必须是void
  • 不能调用C++标准库(除了少数数学函数)
  • 不能使用静态变量

来看一个完整的向量加法例子:

// Kernel定义
__global__ void vecAdd(float* A, 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() {
    // ... 分配内存、拷贝数据 ...
    
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // ... 拷贝结果回CPU ...
}

注意那个if (i < N)的边界检查。数据量不一定正好是Block大小的整数倍,多出来的线程要跳过。我见过有人忘了加这个判断,结果数组越界,程序直接崩了。

核心原则:每个线程只处理一个数据元素,通过索引来区分。这是CUDA编程的基本模式——SIMT(单指令多线程)。

嗯,这里还要提一下性能调优的小技巧。我个人习惯在Kernel里尽量多用局部变量,少用全局变量。因为局部变量会分配到寄存器里,访问速度比Global Memory快两个数量级。

比如这样:

__global__ void betterKernel(float* data, int N) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 局部变量,存在寄存器里
    float local_val = data[tid];
    
    // 做一堆计算
    local_val = local_val * 2.0f + 1.0f;
    local_val = sqrtf(local_val);
    
    // 写回
    data[tid] = local_val;
}

你看,local_val全程在寄存器里,读写都是1个周期。如果每次计算都去读data[tid],那每次都要等400个周期。

最后说一句:写Kernel函数,脑子里要时刻想着「我这个线程在干什么」。每个线程都是独立的,但合起来就是一个并行大军。把任务拆得足够细,让每个线程都有活干,别让GPU闲着。

CUDA编程核心概念结构图 Grid(网格) Block (0,0) T0 T1 T2 ... Shared Memory Block (1,0) T0 T1 T2 ... Shared Memory Block (2,0) T0 T1 T2 ... Shared Memory Global Memory(全局内存) 数据块 A 数据块 B 数据块 C ... Block(线程块) Thread(线程) Global Memory Grid

一句话总结:线程层次结构决定了「谁在干活」,内存层次结构决定了「干活快不快」。两者配合好,你的Kernel才能跑出好性能。