4、设备端内核声明：__global__函数在设备端的调用限制，__device__修饰符的作用

好，咱们今天聊一个挺关键的话题——设备端内核声明。

很多同学刚开始写CUDA的时候，觉得__global__和__device__不就是两个修饰符嘛，加在函数前面就完事了。嗯，我当年也是这么想的。直到有一次我在设备端试图调用一个__global__函数，编译器直接给我甩了一脸错误。那时候我才意识到，这两个修饰符背后藏着不少门道。

4.1 __global__函数：只能在主机端调用

先说说__global__。它的定义很明确：这个函数是内核函数，由主机端（CPU）发起调用，在设备端（GPU）上执行。但有一点非常关键——你不能在设备端代码里直接调用另一个__global__函数。

为什么会这样？

你想想看，__global__函数本质上是一个入口点，它对应着一整个线程网格的启动。如果你在设备端再调用一个__global__，那就意味着要从GPU内部再启动一个新的网格。这在早期的CUDA架构里是不允许的。直到Kepler架构引入了动态并行（Dynamic Parallelism），才允许在设备端通过cudaLaunchDevice或<<<grid, block>>>语法来启动子内核。

但注意，即使有了动态并行，__global__函数本身仍然不能像普通函数那样被直接调用。你得通过特殊的启动语法。说白了，它还是那个「入口点」的角色，只不过现在入口点可以在设备端了。

核心限制：
  • __global__函数不能从设备端直接调用（函数名+参数列表的方式）
  • 在动态并行中，设备端启动__global__必须使用<<<grid, block>>>语法
  • __global__函数必须返回void
  • __global__函数不能是类的成员函数（静态成员函数也不行）
// 错误示例：设备端直接调用 __global__ 函数
__global__ void child_kernel() {
    // 子内核逻辑
}

__global__ void parent_kernel() {
    // 这样写是错的！
    child_kernel();  // 编译错误：不允许从设备端直接调用 __global__ 函数
}

// 正确示例：使用动态并行语法
__global__ void parent_kernel() {
    // 这样写才对
    child_kernel<<<1, 256>>>();  // 动态并行启动子内核
}

我个人习惯在写动态并行代码时，会先检查计算能力是否>=3.5。因为动态并行需要Kepler架构以上的支持。我曾经在项目里吃过这个亏——在Tesla K80上跑得好好的，换到一台老旧的GTX 680上直接崩了。嗯，从那以后我每次都会在代码里加个cudaDeviceGetAttribute检查一下。

4.2 __device__修饰符：设备端函数的正确打开方式

那如果我想在设备端定义一个可以被其他设备函数调用的函数，该怎么办？答案就是__device__修饰符。

__device__函数只能在设备端被调用，而且只能被其他设备函数（__global__或__device__）调用。它不能从主机端直接调用。这个限制其实很合理——你想想看，一个运行在GPU上的函数，CPU怎么直接调用它？

小技巧： 我经常把__device__函数当作「GPU上的工具函数」来用。比如矩阵乘法中的分块计算、归约操作中的warp级操作，都可以封装成__device__函数，让代码更清晰。 
// 定义一个设备端工具函数
__device__ float warp_reduce_sum(float val) {
    // 使用warp shuffle进行归约
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
    }
    return val;
}

// 在内核中调用
__global__ void reduce_kernel(float* input, float* output, int n) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    float sum = 0.0f;
    
    if (tid < n) {
        sum = input[tid];
    }
    
    // 调用 __device__ 函数
    sum = warp_reduce_sum(sum);
    
    // ... 后续处理
}

4.3 __global__与__device__的混合使用

在实际项目中，我经常看到这样的模式：一个__global__内核函数调用多个__device__辅助函数。这种设计非常合理，因为它把复杂的计算逻辑拆解成可复用的模块。

修饰符 调用位置 执行位置 能否被主机调用 能否被设备调用
__global__ 主机端（或动态并行中的设备端） 设备端 是（通过<<<...>>>） 是（通过<<<...>>>，需动态并行）
__device__ 设备端 设备端 否 是（直接调用）
__host__ __device__ 主机端或设备端 主机端或设备端 是 是
避坑指南： 我曾经在写一个双精度矩阵运算时，把一个__device__函数误写成了__global__。结果编译通过了，但运行时每个线程都启动了一个独立的网格，性能直接掉到谷底。调试了半天才发现问题。所以记住：__global__是「启动一个网格」，__device__是「定义一个函数」，两者完全不是一回事。

4.4 动态并行中的递归调用

说到递归，这是动态并行的一个高级用法。你可以在设备端启动一个内核，然后在这个内核里再启动自己——这就形成了递归。

但这里有个坑：递归深度不能太深。因为每次启动子内核都需要分配资源，包括线程块调度、共享内存等。我记得有一次我写了一个二叉树遍历的递归内核，深度到了20层，结果直接爆了显存。后来我限制递归深度不超过8层，才稳定运行。

// 递归内核示例：计算阶乘（仅作演示，实际不推荐这么写）
__global__ void factorial_kernel(int n, long long* result) {
    if (n <= 1) {
        *result = 1;
        return;
    }
    
    long long sub_result;
    // 启动子内核计算 (n-1)!
    factorial_kernel<<<1, 1>>>(n - 1, &sub_result);
    
    // 同步子内核
    cudaDeviceSynchronize();
    
    *result = n * sub_result;
}
递归调用的关键限制：
  • 递归深度受限于设备端的资源（寄存器、共享内存、栈空间）
  • 每次递归调用都需要cudaDeviceSynchronize()来保证子内核完成
  • 递归深度过大会导致性能急剧下降，甚至程序崩溃
  • 建议递归深度不超过8-10层

4.5 我的个人建议

说了这么多，总结几条实战经验吧：

  1. 能用__device__就别用动态并行。动态并行虽然强大，但开销不小。每次启动子内核都有延迟，能在一个内核里用__device__函数解决的问题，就别拆成多个内核。
  2. 检查计算能力。动态并行需要3.5以上的计算能力。我在项目里习惯在初始化时加个断言，避免在旧硬件上跑出奇怪的问题。
  3. 递归深度要控制。如果非要用递归，建议深度不超过8层。更深的话，考虑用迭代或者栈模拟。
  4. 注意同步。动态并行中的子内核启动是异步的，别忘了加cudaDeviceSynchronize()。我曾经漏掉这个同步，结果父内核读取了子内核还没写完的数据，调试了一整天。

好了，关于__global__和__device__在设备端的调用限制，就聊这么多。记住一句话：__global__是启动入口，__device__是工具函数。搞清楚了这一点，你在写CUDA代码时就不会搞混了。

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