3、CUDA编程入门：Kernel函数编写、线程层次结构（Grid/Block/Thread）、<<<>>>语法

好，咱们直接进入正题。CUDA编程，说白了就是让GPU帮你干活。但GPU不是CPU，你不能直接写个for循环扔给它。你得写一种特殊的函数——Kernel函数。然后通过一种特殊的语法——<<<>>>——告诉GPU：嘿，用这么多线程来跑这个函数。

我第一次接触CUDA时，最困惑的就是这个线程模型。CPU上我们习惯了一个线程跑到底，GPU上却是成千上万个线程同时跑。嗯，这里要注意，理解了这个模型，你就掌握了CUDA的精髓。

3.1 Kernel函数：GPU上跑的函数

Kernel函数，就是跑在GPU设备上的函数。它有几个硬性规定：

• 必须用__global__修饰
• 返回类型必须是void
• 不能是类的成员函数（静态成员可以）
• 调用时要用<<>>语法

来看个最简单的例子：

// 定义一个Kernel函数
__global__ void hello_from_gpu() {
    printf("Hello from thread %d in block %d\n", 
           threadIdx.x, blockIdx.x);
}

int main() {
    // 调用Kernel函数
    hello_from_gpu<<<2, 4>>>();
    
    // 同步等待GPU执行完成
    cudaDeviceSynchronize();
    return 0;
}

这个代码会输出什么？2个block，每个block有4个thread，总共8个线程。每个线程都会打印一句话。我刚开始写时，以为它会按顺序打印，结果每次运行顺序都不一样——这就是并行，线程执行顺序是不确定的。

3.2 线程层次结构：Grid、Block、Thread

CUDA的线程组织方式，有点像军队编制：

• Thread（士兵）：最基础的执行单元
• Block（班）：一组线程，可以协作（共享内存、同步）
• Grid（排）：一组Block，构成整个任务

每个线程都有唯一的ID，通过内置变量获取：

变量	含义	示例值
threadIdx.x	线程在block内的索引	0, 1, 2, ..., blockDim.x-1
blockIdx.x	block在grid内的索引	0, 1, 2, ..., gridDim.x-1
blockDim.x	每个block的线程数	由<<<>>>第二个参数决定
gridDim.x	grid内的block数	由<<<>>>第一个参数决定

要计算全局唯一的线程ID，公式是：

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

这个公式我建议你背下来。我在项目中调试过无数次，发现很多bug都是因为ID算错了。

3.3 <<<>>>语法：启动配置

这个语法看起来有点怪，但用多了就习惯了。它的完整形式是：

kernel_name<<<grid_dim, block_dim, shared_mem_bytes, stream>>>(args...);

四个参数：

1. grid_dim：grid的维度，可以是int或dim3类型
2. block_dim：block的维度，可以是int或dim3类型
3. shared_mem_bytes：动态共享内存大小（可选，默认0）
4. stream：CUDA流（可选，默认0）

一维情况最简单：

// 启动2个block，每个block4个线程
my_kernel<<<2, 4>>>();

二维或三维情况用dim3：

// 2D grid: 2x3 = 6个block
// 2D block: 4x2 = 8个线程/block
dim3 grid_dim(2, 3);
dim3 block_dim(4, 2);
my_kernel<<<grid_dim, block_dim>>>();

为什么会需要多维？因为图像、矩阵这些数据天然就是二维的。用多维索引，代码更直观。

3.4 实战：向量加法

来写个完整的例子。向量加法，就是把两个数组对应元素相加：

#include <cuda_runtime.h>
#include <stdio.h>

// Kernel函数：向量加法
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查：防止越界
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

int main() {
    int n = 1000;
    size_t bytes = n * sizeof(float);
    
    // 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);
    
    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    // 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    
    // 启动Kernel
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 拷贝结果回CPU
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }
    
    // 释放内存
    free(h_a); free(h_b); free(h_c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    
    return 0;
}

注意那个blocks_per_grid的计算：(n + threads_per_block - 1) / threads_per_block。这是向上取整的技巧。比如1000个元素，256线程/block，需要4个block（前3个满的，最后一个只处理232个）。

3.5 线程层次结构图

下面这张图，我画了线程层次结构的核心关系。你看一眼就明白了：

这张图展示的是2D的线程层次。实际项目中，一维用得最多，但理解多维有助于处理图像、矩阵这类数据。

3.6 常见坑与避坑指南

嗯，这一章的内容就到这。CUDA编程入门，说白了就是三件事：写Kernel、算ID、配<<<>>>。把这三点搞明白了，后面优化显存、调优性能才有基础。