2. Hello World核函数：从零开始写你的第一个CUDA程序

说实话，我教了这么多年CUDA，发现很多新手都卡在第一步——不是语法看不懂，而是搞不清CPU和GPU到底怎么配合工作的。这一章，我们就从最基础的核函数语法讲起，手把手带你写出第一个向量加法。

2.1 核函数定义语法：给GPU下指令

核函数，说白了就是在GPU上执行的函数。它跟普通的C函数长得有点像，但有几个关键区别。

先看一个最简单的例子：

__global__ void hello_kernel() {
    printf("Hello from thread %d\n", threadIdx.x);
}

注意那个 __global__ 修饰符。这是CUDA的“身份证”，告诉编译器：这个函数要跑在GPU上，由CPU来调用。我个人习惯把 __global__ 放在函数返回类型前面，这样一眼就能看出来是核函数。

核函数有几个硬性规定：

  • 必须返回 void 类型
  • 不能是类的成员函数（静态成员函数可以）
  • 调用时要用 <<<...>>> 语法指定执行配置
  • 函数体内可以访问 threadIdx、blockIdx 等内置变量

嗯，这里要注意：__global__ 核函数不能直接递归调用，也不能有可变参数。我在项目中遇到过有人试图传 ... 参数进去，结果编译报错半天没找到原因。

2.2 <<<执行配置>>>：告诉GPU开多少线程

调用核函数时，<<<gridDim, blockDim>>> 这部分就是执行配置。它决定了GPU要开多少个线程来干活。

hello_kernel<<<1, 10>>>();

这行代码的意思是：启动1个线程块（block），每个块里有10个线程。总共就是10个线程在跑。

你想想看，如果改成 <<<2, 5>>>，那就是2个块，每个块5个线程，总共还是10个线程。但线程的组织方式变了。

为什么会这样设计？因为GPU的硬件结构就是按“块”来组织的。一个块内的线程可以共享数据，块之间是独立的。我刚开始学的时候，总把 gridDim 和 blockDim 搞反，后来记住一句话：grid是块的集合，block是线程的集合。

核心概念：gridDim 和 blockDim 可以是三维的，比如 <<<dim3(2,3,1), dim3(4,4,1)>>>。但大多数情况下，我们用一维就够了。

2.3 主机与设备内存模型：CPU和GPU各管各的

这是新手最容易踩坑的地方。CPU（主机）和GPU（设备）有各自独立的内存空间。你不能直接在GPU上访问CPU的变量，反过来也不行。

看这张图就明白了：

主机 (Host) CPU 内存 (DDR) 变量 a, b, c cudaMemcpyHostToDevice cudaMemcpyDeviceToHost 设备 (Device) GPU 显存 (VRAM) d_a, d_b, d_c

数据传递的流程是固定的：

  1. 在主机端分配内存（malloc 或 new）
  2. 在设备端分配显存（cudaMalloc）
  3. 用 cudaMemcpy 把数据从主机拷贝到设备
  4. 启动核函数在GPU上计算
  5. 用 cudaMemcpy 把结果从设备拷回主机
  6. 释放两边的内存

我曾经犯过一个低级错误：在核函数里直接访问主机指针。结果程序崩溃，查了半天才发现是内存空间不对。记住：核函数里只能访问设备指针，主机指针传进去也没用。

2.4 第一个向量加法实现：动手写代码

好了，理论讲完了，我们来写一个完整的向量加法。假设有两个长度为N的浮点数组A和B，我们要计算 C = A + B。

先看完整代码：

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

#define N 1024

// 核函数：每个线程计算一个元素
__global__ void vec_add(float *d_a, float *d_b, float *d_c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        d_c[idx] = d_a[idx] + d_b[idx];
    }
}

int main() {
    float *h_a, *h_b, *h_c;          // 主机指针
    float *d_a, *d_b, *d_c;          // 设备指针
    
    // 1. 主机分配内存
    h_a = (float*)malloc(N * sizeof(float));
    h_b = (float*)malloc(N * sizeof(float));
    h_c = (float*)malloc(N * sizeof(float));
    
    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }
    
    // 2. 设备分配显存
    cudaMalloc(&d_a, N * sizeof(float));
    cudaMalloc(&d_b, N * sizeof(float));
    cudaMalloc(&d_c, N * sizeof(float));
    
    // 3. 数据从主机拷贝到设备
    cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);
    
    // 4. 启动核函数
    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);
    
    // 5. 结果拷回主机
    cudaMemcpy(h_c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("h_c[%d] = %f\n", i, h_c[i]);
    }
    
    // 6. 释放内存
    free(h_a); free(h_b); free(h_c);
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    
    return 0;
}

这段代码有几个关键点，我挨个说一下：

  • 线程索引计算：idx = threadIdx.x + blockIdx.x * blockDim.x 是标准写法。每个线程通过这个公式拿到自己负责的元素下标。
  • 边界检查：if (idx < n) 这行不能省。因为线程总数可能比N大，多余的线程什么都不做就行。
  • 线程块数量计算：(N + threads_per_block - 1) / threads_per_block 是向上取整的技巧。比如N=1024，每块256线程，那就是4个块。

我的习惯：线程块大小一般选128或256。太小了硬件利用率低，太大了寄存器不够用。256是个比较稳妥的起点。

2.5 常见错误与调试技巧

写CUDA程序，最烦人的就是错误不好查。我分享几个经验：

  • 检查返回值：每个CUDA API调用都应该检查返回值。可以用 cudaGetLastError() 捕获核函数错误。
  • 用 cuda-memcheck：这个工具能帮你检查越界访问、未初始化内存等问题。我每次调试必开。
  • 先跑小数据：别一上来就搞几百万个元素。先设N=10，打印结果看看对不对。

说实话，向量加法虽然简单，但它包含了CUDA编程的所有核心要素：内存管理、线程组织、数据搬运。把这个搞懂了，后面学矩阵乘法、卷积这些复杂算子就会轻松很多。

好了，这一章就到这里。代码我已经跑过很多次了，你直接复制就能用。如果遇到问题，多半是环境配置或者指针搞混了——别急，慢慢排查。

可可精品资料库

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