第4章：核函数编写基础：__global__、__device__、__host__修饰符、核函数参数传递

好，咱们今天聊聊核函数编写的基础。说实话，很多新手一上来就急着写复杂的并行算法，结果连修饰符都用错了，调试起来那叫一个痛苦。我个人习惯是，先把这些基础概念吃透，后面写代码才能行云流水。

4.1 三个修饰符：各司其职

CUDA里有三个核心修饰符：__global__、__device__、__host__。它们决定了函数在哪儿跑、谁来调用。

修饰符	执行位置	调用位置	典型用途
__global__	GPU	CPU（主机端）	核函数入口
__device__	GPU	GPU（设备端）	被核函数调用的辅助函数
__host__	CPU	CPU（主机端）	普通CPU函数（默认）

__global__ 是核函数的标志。你想想看，它必须返回 void，而且只能从CPU端调用。我在项目中遇到过有人试图在核函数里再调用另一个 __global__ 函数，编译直接报错——这是不允许的。

__device__ 函数只能在GPU上被其他GPU函数调用。说白了，它就是给核函数当“小弟”的。比如你要在核函数里算个复杂的数学公式，可以拆成几个 __device__ 函数，代码更清晰。

__host__ 是默认的，不写也是它。但有个技巧：你可以同时写 __host__ __device__，这样同一个函数既能在CPU跑，也能在GPU跑。嗯，这里要注意，这种函数里不能有 printf 之类的设备端不支持的操作。

4.2 核函数参数传递：没那么简单

核函数的参数传递，我刚开始学的时候也踩过坑。你以为跟C语言一样传个指针就完事了？没那么简单。

先看个标准例子：

__global__ void vecAdd(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];
    }
}

// 调用方式
vecAdd<<<grid, block>>>(d_A, d_B, d_C, N);

这里有几个关键点：

• 参数默认按值传递：所有参数都会被复制到GPU的常量内存中。对于小数据（比如 int N），这没问题。但如果你传一个巨大的结构体，那就炸了——常量内存只有64KB。
• 指针传递的是设备端地址：d_A 必须是 cudaMalloc 分配的设备指针。我曾经见过有人把CPU端的指针传进去，结果核函数里访问的全是垃圾数据。
• 参数数量有限制：我记得好像是最多255个参数？但实际没人会传那么多。超过4-5个参数时，我建议用结构体打包。

4.3 参数传递的几种模式

根据我的经验，参数传递有三种常见模式：

模式	适用场景	示例
直接传值	标量参数（int, float等）	vecAdd<<<...>>>(d_A, d_B, d_C, N);
传指针	数组、大块数据	d_A 是设备指针
传结构体	多个相关参数打包	Params p = {d_A, d_B, d_C, N};

我个人比较推荐第三种。当参数超过3个时，用结构体打包不仅代码整洁，还能减少寄存器压力。你想想看，每个线程都要从常量内存读参数，如果参数太多，寄存器都不够用了。

4.4 知识体系总览

下面这张图是我自己画的，把本章的核心逻辑串起来了。你看一眼就能明白修饰符和参数传递的关系：

这张图把修饰符和参数传递的关系讲得很清楚。你从上往下看：先选对修饰符，再决定怎么传参，最后注意那些容易踩的坑。

4.5 实战建议

最后，我分享几个实战中的小建议：

• 能用 __device__ 就别往核函数里塞太多逻辑：核函数本身应该尽量简洁，复杂计算拆成多个 __device__ 函数，方便调试和复用。
• 参数传递优先用结构体：当参数超过3个时，结构体能减少寄存器压力，还能让代码更易读。
• 检查指针合法性：传参前用 cudaPointerAttributes 确认指针是设备端的，别把主机指针传进去。
• 常量内存别滥用：每个核函数的参数都会占用常量内存，如果参数太多或太大，考虑用 cudaMemcpyToSymbol 提前存好。