__global__ && __device__

global && device

内容

在CUDA编程中，__global__和__device__是用于定义GPU函数的关键字，但它们的用途和调用方式有本质区别。以下是两者的详细对比：

---

核心区别

特性	__global__ 函数	__device__ 函数
调用者	由CPU（主机）调用	只能由GPU代码（其他__device__或__global__函数）调用
执行位置	在GPU上执行（称为内核函数）	在GPU上执行（称为设备函数）
返回值	必须返回void	可以有任意返回值
调用语法	使用<<<grid, block>>>指定线程配置	像普通函数一样直接调用
典型用途	主计算逻辑（并行任务入口）	辅助计算（数学运算、数据预处理等）

---

关键细节

1. __global__ 函数的限制

• 参数传递：只能通过值传递（不能传递主机内存指针）

  __global__ void kernel(int a, float* dev_ptr) { ... }  // 正确
  __global__ void kernel(int& a) { ... }                  // 错误（引用传递）
  

• 动态并行：需要计算能力≥3.5的GPU才支持从设备端调用__global__函数

  __device__ void dev_func() {
      kernel<<<1,1>>>();  // 需要启用动态并行编译选项
  }
  

2. __device__ 函数的扩展特性

• 与__host__联合使用：实现CPU/GPU通用函数

  __host__ __device__ int add(int a, int b) {
      return a + b;  // 主机和设备代码均可调用
  }
  

• 内存访问：可直接操作全局内存（cudaMalloc分配的指针）

  __device__ void modify(float* data) {
      data[threadIdx.x] *= 2.0f;  // 直接修改全局内存
  }
  

---

代码示例

场景：向量加法（__global__调用__device__）

// 设备函数（GPU内部辅助计算）
__device__ float element_add(float a, float b) {
    return a + b;
}

// 全局函数（内核入口）
__global__ void vector_add(float* x, float* y, float* z, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        z[idx] = element_add(x[idx], y[idx]);  // 调用设备函数
    }
}

// 主机调用
int main() {
    float *d_x, *d_y, *d_z;
    cudaMalloc(&d_x, size);  // 分配设备内存
    cudaMalloc(&d_y, size);
    cudaMalloc(&d_z, size);

    // 启动内核
    vector_add<<<grid, block>>>(d_x, d_y, d_z, N);
}

---

性能注意事项

1. __device__内联优化：

   • 使用__forceinline__强制内联小型设备函数，避免函数调用开销

   __device__ __forceinline__ float fast_sqrt(float x) {
       return sqrtf(x);  // 内联后减少上下文切换
   }
   

2. 寄存器压力：

   • 复杂的__device__函数可能导致寄存器溢出，使用launch_bounds优化

   __global__ void __launch_bounds__(256, 4) my_kernel() { ... }
   

---

常见误区

1. 错误尝试从主机调用__device__函数：

   __device__ void dev_func() { ... }
   
   int main() {
       dev_func();  // 编译错误：主机无法调用设备函数
   }
   

2. 误用__global__返回值：

   __global__ int kernel() { return 1; }  // 错误：必须返回void
   

---

总结

• __global__：GPU计算的入口点，由CPU触发大规模并行任务
• __device__：GPU内部的工具函数，用于代码复用和模块化设计
• 组合使用：通过__global__组织并行架构，用__device__封装可复用的计算单元