第三章：CUDA编程核心概念

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

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

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

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

• Thread（线程）：最基础的执行单元
• Block（线程块）：一组线程，可以协作
• Grid（网格）：一组Block，构成整个任务

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

我在项目中遇到过一个问题：有个同事把数据分成了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;

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是你手里的工具。你总不能每次拧螺丝都跑一趟仓库拿扳手吧？

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大小的整数倍，多出来的线程要跳过。我见过有人忘了加这个判断，结果数组越界，程序直接崩了。

嗯，这里还要提一下性能调优的小技巧。我个人习惯在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闲着。