第一个CUDA程序：核函数编写、线程层次结构、<<<>>>语法详解、编译与运行

好，咱们直接上手写代码。很多初学者觉得CUDA门槛高，其实说白了，第一个程序比你想的简单得多。我当年刚接触CUDA时，花了整整一下午才把环境配好，结果第一个"Hello World"跑了三小时才跑通——后来发现是线程索引算错了。嗯，这些坑我今天帮你一次性避开。

1. 核函数：GPU上执行的函数

核函数（Kernel Function）就是在GPU上跑的函数。它有几个硬性规定：

  • 必须用 __global__ 修饰
  • 返回类型必须是 void
  • 不能使用可变参数（比如 ...）
  • 不能是类的成员函数（静态成员函数可以）

来看一个最简单的例子：

// 定义一个核函数
__global__ void hello_from_gpu() {
    printf("Hello from GPU! Thread ID: %d\n", threadIdx.x);
}

注意那个 threadIdx.x，这是CUDA内置变量，代表当前线程在块内的索引。我刚开始写的时候，总把它和 blockIdx.x 搞混——这两个东西后面会细讲。

个人习惯：我写核函数时，函数名都会加个 _kernel 后缀，比如 vec_add_kernel。这样在代码里一眼就能看出哪些是GPU函数，哪些是CPU函数。

2. 线程层次结构：网格、块、线程

CUDA的线程组织方式，说白了就是三层嵌套：

  • 网格（Grid）：最外层，包含多个线程块
  • 线程块（Block）：中间层，包含多个线程
  • 线程（Thread）：最内层，实际执行计算的最小单位

你可以想象成：一个网格是一个大工厂，每个线程块是一个车间，每个线程是一个工人。工人之间可以协作（通过共享内存），但不同车间的工人不能直接通信。

关键点：每个线程都有唯一的ID。这个ID由 blockIdx 和 threadIdx 共同决定。全局线程ID的计算公式是： 
int tid = blockIdx.x * blockDim.x + threadIdx.x;
这个公式我建议你背下来——几乎每个CUDA程序都会用到。

为什么会这样设计？你想想看，如果所有线程都用一个全局ID，那管理起来多麻烦。分块之后，每个块内的线程可以共享数据，效率高得多。我在项目中遇到过，有人把所有线程都塞到一个块里，结果性能惨不忍睹——因为一个块最多只能有1024个线程（现代GPU）。

3. <<<>>> 语法详解：启动配置

这个语法可能是CUDA里最"怪异"的地方了。它用来指定启动多少个线程：

// 语法：kernel_name<<<grid_dim, block_dim>>>(args);
hello_from_gpu<<<1, 10>>>();

这里 <<<1, 10>>> 的意思是：启动1个线程块，每个块里有10个线程。总共就是10个线程。

参数可以是整数，也可以是 dim3 类型：

// 一维网格，一维块
hello_from_gpu<<<2, 5>>>();

// 二维网格，二维块（更灵活）
dim3 grid_dim(2, 3);   // 2x3 = 6个块
dim3 block_dim(4, 4);  // 4x4 = 16个线程/块
hello_from_gpu<<<grid_dim, block_dim>>>();

我记得第一次用二维网格时，死活想不通为什么 threadIdx.y 是行索引而不是列索引。后来才明白，CUDA用的是行主序——跟C语言一样。

避坑指南：我曾经在项目里把 blockDim.x 和 gridDim.x 搞反了，结果数据全算错了。记住：blockDim 是每个块的大小，gridDim 是网格里有多少个块。别问我怎么知道的...

4. 完整的CUDA程序：向量加法

理论说完了，咱们写个完整的例子。向量加法是CUDA的"Hello World"：

#include <stdio.h>

// 核函数：向量加法
__global__ void vec_add_kernel(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 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. 拷贝数据到GPU
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    
    // 4. 启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    
    vec_add_kernel<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);
    
    // 5. 拷贝结果回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;
}

这个流程你记一下：分配内存 → 拷贝到GPU → 启动核函数 → 拷贝回CPU → 清理。几乎所有CUDA程序都是这个套路。

个人经验：那个 (n + threads_per_block - 1) / threads_per_block 是向上取整的经典写法。我习惯用256作为每个块的线程数，因为这个值对大多数GPU来说性能都不错。太小了浪费硬件资源，太大了可能超过块的最大线程数限制。

5. 编译与运行

编译CUDA程序用 nvcc 编译器：

# 编译
nvcc -o vec_add vec_add.cu

# 运行
./vec_add

如果你有多个文件，可以这样：

nvcc -c kernel.cu -o kernel.o
nvcc -c main.cpp -o main.o
nvcc kernel.o main.o -o program

嗯，这里要注意：.cu 文件是CUDA源文件，nvcc 会自动识别并调用GPU编译器。如果你混用C++代码，记得用 -x cu 选项告诉编译器。

避坑指南：我曾经在编译时忘了加 -arch=sm_75（对应我的GPU架构），结果程序跑起来慢得像蜗牛。建议你用 nvcc --list-gpu-arch 查看你的GPU支持哪些架构，然后指定对应的 -arch 参数。

6. 知识体系总览

下面这张图帮你理清本章的核心逻辑：

第一个CUDA程序：核心知识体系 核函数 __global__ void func() 线程层次结构 Grid → Block → Thread <<<>>> 语法 启动配置参数 返回void __global__修饰 不能可变参数 静态成员可 threadIdx 块内索引 blockIdx 网格内索引 blockDim 块大小 grid_dim 网格维度 block_dim 块维度 编译与运行 nvcc -o program source.cu 完整程序流程 分配内存 → cudaMemcpy到GPU → 启动核函数 → cudaMemcpy回CPU → 清理

7. 调试小技巧

写CUDA程序最头疼的就是调试。我分享几个实用技巧：

  • 检查错误码：每次调用CUDA API后，用 cudaGetLastError() 检查是否有错误。我习惯写个宏：#define CUDA_CHECK(call) { cudaError_t err = call; if (err != cudaSuccess) ... }
  • 用 cuda-memcheck：这个工具能帮你检测越界访问、未初始化内存等问题。运行方式：cuda-memcheck ./program
  • 打印调试信息：核函数里可以用 printf，但注意不要打印太多，否则会拖慢性能。我一般只在调试时用，正式代码里会注释掉。
个人经验：我刚开始写CUDA时，经常遇到"核函数没执行"的问题。后来发现是忘了加 cudaDeviceSynchronize()——这个函数会等待所有GPU任务完成。记住：CPU和GPU是异步的，不加同步的话，CPU可能已经跑到下一行了，GPU还没开始干活。

好了，第一个CUDA程序的核心内容就这些。你照着上面的代码敲一遍，跑通了，就算入门了。别怕出错，我当年第一个程序跑了三小时才跑通，你肯定比我快。

可可精品资料库
← 上一章 📖 返回目录 下一章 →