3、CUDA基础语法：核函数编写、线程层次结构、内存模型、第一个CUDA程序

好，咱们正式开始写代码了。

前面聊了那么多GPU的硬件架构，说白了，都是为了这一刻做铺垫。你想想看，如果连线程怎么组织、数据放哪儿都不知道，那再强的Tensor Core也跟你没关系。

这一章，我带你亲手写一个CUDA程序。别怕，代码量不大，但背后的概念——核函数、线程层次、内存模型——是后续所有高性能计算的地基。我自己带团队时，新人最容易在这几个地方栽跟头，咱们今天一次性讲透。

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

核函数是什么？简单说，就是在GPU上执行的函数。CPU调用它，GPU执行它。

写法上有个硬性规定：前面必须加 __global__ 修饰符。返回值必须是 void。我刚开始学的时候，老记不住这个 void，编译报错半天才发现——嗯，这种低级错误我犯过不止一次。

核函数定义三要素：

  • 修饰符：__global__（告诉编译器这是给GPU跑的）
  • 返回值：void（不能返回任何值，结果通过参数传回来）
  • 调用方式：使用三括号语法 <<<grid, block>>>

看个最简单的例子：

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

// CPU端调用
int main() {
    float *d_data;
    int n = 1024;
    // ... 分配显存、拷贝数据 ...
    add_one<<<1, 256>>>(d_data, n);  // 1个block，256个线程
    // ... 拷贝回结果 ...
}

这里有个细节：核函数里用 threadIdx.x、blockIdx.x、blockDim.x 来算全局索引。为什么这么算？因为GPU的线程是分层次组织的，你得告诉每个线程：你是谁，该处理哪份数据。

避坑指南：

我曾经在项目里忘了加 if (idx < n) 这个边界检查。数据量刚好是线程数的整数倍时没问题，一旦不是，数组越界，结果全乱套。所以——永远、永远、永远要做边界检查。

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

GPU的线程不是一盘散沙，而是有组织有纪律的。层次结构就三层：

  • Thread（线程）：最小编程单元，每个线程执行同一段代码，但处理不同的数据
  • Block（线程块）：一组线程，可以协作（共享内存、同步）
  • Grid（网格）：一组Block，构成整个任务

调用时用 <<<grid_dim, block_dim>>> 指定。比如：

// 启动一个Grid，包含2个Block，每个Block有256个线程
kernel_name<<<2, 256>>>(args);

// 也可以多维
dim3 grid_dim(4, 4);    // 4x4 = 16个Block
dim3 block_dim(16, 16); // 16x16 = 256个线程
kernel_name<<<grid_dim, block_dim>>>(args);

为什么要有Block？说白了，是为了让线程能协作。同一个Block里的线程可以访问共享内存，可以同步。不同Block之间，老死不相往来——这是硬件决定的。

我个人的习惯是：Block大小选128或256，通常是32的倍数（因为warp大小是32）。选太小，硬件利用率低；选太大，寄存器不够用，性能反而下降。这个平衡点，得靠调参试出来。

关键内置变量：

变量含义
threadIdx.x当前线程在Block内的索引
blockIdx.x当前Block在Grid内的索引
blockDim.x每个Block的线程数
gridDim.xGrid内的Block数

算全局索引的公式：

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

这个公式你得刻在脑子里。我面试新人时，第一个问题就是这个。

3.3 内存模型——数据在哪儿，性能就在哪儿

GPU的内存不是一锅粥，分好几层。每一层的速度、容量、作用都不一样。你想想看，如果所有数据都从显存读，那带宽再大也扛不住。

主要的内存类型：

内存类型位置速度作用域生命周期
全局内存显存慢（~200GB/s）所有线程程序结束
共享内存芯片内快（~1TB/s）同一个BlockBlock结束
寄存器SM内极快单个线程线程结束
常量内存显存（有缓存）快（只读）所有线程程序结束

我在项目中遇到过一个问题：核函数里频繁访问全局内存，性能死活上不去。后来把热点数据搬到共享内存里，速度直接翻了3倍。说白了，优化CUDA程序，80%的工作就是在折腾数据放哪儿。

注意：

共享内存虽然快，但容量有限（通常48KB~96KB）。别一股脑全塞进去，会爆的。我见过有人把整个数组拷进共享内存，结果Block启动不了——嗯，这种错误我也犯过。

3.4 第一个完整的CUDA程序

光说不练假把式。咱们写一个完整的程序：两个数组相加。

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

// 核函数：向量加法
__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 1024;
    size_t size = n * sizeof(float);

    // 1. 在CPU上分配内存
    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }

    // 2. 在GPU上分配显存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 3. 把数据从CPU拷贝到GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int block_size = 256;
    int grid_size = (n + block_size - 1) / block_size;
    vec_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);

    // 5. 把结果从GPU拷回CPU
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("c[%d] = %f\n", i, h_c[i]);
    }

    // 7. 释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

这个流程你记一下：

  1. CPU分配内存，初始化数据
  2. GPU分配显存
  3. 数据从CPU拷到GPU
  4. 启动核函数
  5. 结果从GPU拷回CPU
  6. 验证结果
  7. 释放所有内存

少一步都不行。我刚开始写的时候，经常忘了 cudaMemcpy 拷回来，结果CPU上全是垃圾数据——排查了半天才发现。

调试小技巧：

每次调用CUDA API后，加一句 cudaGetLastError() 检查错误。比如：

vec_add<<<grid_size, block_size>>>(d_a, d_b, d_c, n);
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
}

这个习惯能帮你省下大量调试时间。我个人的项目里，所有核函数调用后面都跟着这个检查。

3.5 小结

这一章的内容，说白了就是三件事：

  • 核函数怎么写（__global__ void + 三括号）
  • 线程怎么组织（Grid、Block、Thread）
  • 数据放哪儿（全局内存、共享内存、寄存器）

这些是CUDA编程的骨架。后面的Tensor Core编程、矩阵乘法优化，全是在这个骨架上添肉。你把这章吃透了，后面就顺了。

下一章，咱们开始玩真的——用Tensor Core做矩阵乘法。到时候你会发现，今天学的线程索引计算，会一遍又一遍地用上。

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