第一章：CUDA编程基础——主机与设备、核函数、线程层次结构

大家好，我是你们这门课的老师。在GPU计算领域摸爬滚打了十几年，踩过的坑比写过的代码还多。今天咱们开始第一讲，聊聊CUDA编程最基础的东西。

说实话，很多人一上来就急着写复杂的算子，结果连最基本的“主机”和“设备”都没搞清楚。我见过太多人把CPU和GPU的内存混为一谈，最后调试到崩溃。所以，咱们先把地基打牢。

1.1 主机(Host)与设备(Device)

在CUDA的世界里，主机指的是CPU及其内存系统，设备指的是GPU及其显存。这两个家伙分工明确：

  • 主机(CPU)：擅长逻辑控制、串行任务、少量数据的管理
  • 设备(GPU)：擅长大规模并行计算、大量数据的吞吐

你想想看，CPU就像个精明的项目经理，GPU则是一群任劳任怨的搬砖工。项目经理负责分配任务、检查结果，搬砖工只管埋头干活。

核心要点：主机和设备有各自独立的内存空间。主机内存用 malloc 分配，设备显存用 cudaMalloc 分配。数据不能直接互通，必须通过 cudaMemcpy 显式拷贝。

我在项目中遇到过一位同事，他直接在GPU核函数里访问主机指针，结果程序崩溃了三天都没找到原因。嗯，这就是没搞懂内存隔离的后果。

1.2 核函数(Kernel)的编写

核函数就是在GPU上执行的函数。它的写法有讲究：

// 核函数定义：前面加 __global__ 修饰符
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

注意几个关键点：

  • 必须用 __global__ 修饰，告诉编译器这是给GPU跑的
  • 返回类型必须是 void
  • 不能使用C++的STL容器，不能有变长参数
  • 可以调用设备函数（用 __device__ 修饰）

个人习惯：我写核函数时，总会在参数列表最后加一个 int N 表示数据长度，然后在函数体内做边界检查。这能避免很多越界问题，尤其是当数据量不是线程块大小的整数倍时。

1.3 线程层次结构：Grid, Block, Thread

这是CUDA最核心的概念，也是新手最容易晕的地方。我画了一张图，帮你理清关系：

Grid (网格) Block (0,0) Block (1,0) Block (2,0) T(0,0) T(1,0) T(0,1) T(1,1) ... ... Grid Block Thread

说白了，层级关系是这样的：

  • Thread（线程）：最小的执行单元，每个线程执行核函数的一份拷贝
  • Block（线程块）：一组线程的集合，它们可以协作（通过共享内存、同步操作）
  • Grid（网格）：一组线程块的集合，构成一次核函数启动

每个线程都有一个唯一的ID，通过 threadIdx、blockIdx、blockDim 来计算：

// 一维情况
int tid = threadIdx.x + blockIdx.x * blockDim.x;

// 二维情况（常用在图像处理）
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
int idx = y * width + x;

我曾经踩过的坑：刚开始写CUDA时，我总以为 threadIdx.x 就是全局唯一的线程ID。结果在多block的情况下，每个block的 threadIdx.x 都是从0开始的！必须结合 blockIdx 和 blockDim 才能算出全局索引。这个错误让我debug了整整一个下午。

1.4 <<< >>> 语法详解

这个语法看起来有点怪，但它是CUDA启动核函数的专用方式：

// 语法：kernel_name<<<grid_dim, block_dim>>>(args...);

// 一维示例
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

// 二维示例
dim3 blockDim(16, 16);
dim3 gridDim((width + 15) / 16, (height + 15) / 16);
kernel2D<<<gridDim, blockDim>>>(d_data, width, height);

参数说明：

参数位置 含义 类型 说明
第一个 Grid维度 dim3 或 int 指定有多少个block
第二个 Block维度 dim3 或 int 指定每个block有多少个thread
第三个（可选） 共享内存大小 size_t 每个block动态分配的共享内存字节数
第四个（可选） 流(Stream) cudaStream_t 指定在哪个流上执行

我的建议：新手先用一维的 <<<blocks, threads>>> 就好。等熟悉了再玩二维三维。另外，block内的线程数最好是32的倍数（warp大小），比如128、256、512。我一般用256，兼容性最好。

1.5 完整的向量加法示例

把上面所有知识点串起来，看一个完整的例子：

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

// 核函数
__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = threadIdx.x + blockIdx.x * blockDim.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    int N = 1 << 20;  // 约100万个元素
    size_t size = N * sizeof(float);
    
    // 1. 主机端分配内存
    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. 设备端分配显存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);
    
    // 3. 数据从主机拷贝到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    
    // 4. 启动核函数
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    
    // 5. 结果拷贝回主机
    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;
}

这个流程你记牢了：分配 → 拷贝到设备 → 启动核函数 → 拷贝回主机 → 清理。我刚开始学的时候，每次写代码都按这个步骤来，慢慢就形成肌肉记忆了。

避坑指南：别忘了检查cuda API的返回值！用 cudaGetLastError() 或 cudaError_t 捕获错误。我曾经因为忘记检查 cudaMalloc 是否成功，结果在显存不足时程序静默崩溃，查了两天才发现是显存爆了。

好了，第一章的内容就到这里。记住，CUDA编程的核心就是理解主机与设备的协作关系，掌握线程层次结构，熟练使用 <<< >>> 语法。这些基础打牢了，后面学算子优化才会事半功倍。

可可精品资料库

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