4. CUDA C基础语法：函数修饰符、变量修饰符与内置变量

各位同学，今天我们来聊聊CUDA C里最基础、也最绕不开的几个概念。说白了，就是告诉GPU：这段代码谁来跑？数据放哪里？我在哪个位置？

我刚开始接触CUDA时，被__global__、__device__、__host__这几个修饰符搞得有点晕。后来在项目中写了一个核函数，忘了加__global__，编译报错找了半小时……嗯，从那以后我就再也不敢小看这些修饰符了。

4.1 函数修饰符：告诉GPU谁干活

CUDA里，函数修饰符决定了这个函数在哪个设备上执行，以及从哪里调用。一共就三种，咱们一个一个说。

修饰符	执行位置	调用位置	典型用途
__global__	GPU（设备端）	CPU（主机端）	核函数入口，启动并行线程
__device__	GPU（设备端）	GPU（设备端）	被核函数调用的辅助函数
__host__	CPU（主机端）	CPU（主机端）	普通CPU函数（可省略）

__global__：核函数的入口

这个修饰符最常用。你写的并行计算代码，入口函数必须是__global__。它返回类型必须是void，不能有返回值。

// 一个简单的向量加法核函数
__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>>>(A, B, C, N);。我个人习惯把网格和块的大小定义成变量，方便调试时改。

__device__：GPU上的小帮手

有时候核函数里需要调用一些辅助计算，比如求平方根、做矩阵小运算。这些函数用__device__修饰，只能在GPU上被其他GPU函数调用。

// 一个设备端辅助函数
__device__ float square(float x) {
    return x * x;
}

__global__ void compute(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] = square(data[idx]);  // 调用设备函数
    }
}

你想想看，如果把这个函数写成普通CPU函数，核函数里调用它，编译会报错——因为GPU不认识CPU函数。反过来也一样。

__host__：默认就是它

不加任何修饰符的函数，默认就是__host__。它只在CPU上跑。你可以显式写__host__，但没必要。不过有一种情况有用：同时定义__host__ __device__，表示这个函数既能在CPU跑，也能在GPU跑。

// 双平台通用函数
__host__ __device__ float clamp(float x, float minVal, float maxVal) {
    return fminf(fmaxf(x, minVal), maxVal);
}

我在做量化回测引擎时，有些数学函数需要在CPU和GPU两端都用，就用这种写法。省得写两遍代码，维护起来也方便。

4.2 变量修饰符：数据放哪里

变量修饰符决定了数据存储在哪个内存区域。CUDA的内存模型比较复杂，但咱们先记住最常用的几个。

修饰符	存储位置	作用域	生命周期
__device__	全局显存	全局	应用程序运行期间
__constant__	常量内存	全局	应用程序运行期间
__shared__	共享内存	线程块内	线程块生命周期

__device__：全局显存变量

声明在全局作用域，所有核函数都能访问。但注意，CPU端不能直接读写它，需要用cudaMemcpy之类的API。

__device__ float global_scale = 1.0f;  // 全局显存变量

__global__ void scaleData(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] *= global_scale;
    }
}

__constant__：只读的常量

常量内存有缓存，读取速度比全局显存快。适合存储那些不会变的参数，比如模型权重、阈值等。我建议把量化策略里的固定参数放这里。

__constant__ float coef[4];  // 64KB上限，注意别超

__global__ void applyModel(float *data, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        data[idx] = coef[0] * data[idx] + coef[1];
    }
}

__shared__：线程块的共享空间

共享内存是每个线程块私有的，块内所有线程都能读写。速度极快，但容量小（通常几十KB）。适合做数据缓存、归约操作。

__global__ void sharedReduce(float *input, float *output, int N) {
    __shared__ float cache[256];  // 每个块有自己的cache
    
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int tid = threadIdx.x;
    
    cache[tid] = (idx < N) ? input[idx] : 0.0f;
    __syncthreads();  // 同步，确保所有线程写完了
    
    // 归约操作...
}

这里有个坑：一定要加__syncthreads()。我曾经在写归约时忘了同步，结果有些线程读到的是旧数据，算出来的结果全是错的。找bug找了半天……

4.3 内置变量：我在哪？

CUDA提供了几个内置变量，让你知道当前线程在网格中的位置。这些变量不需要声明，直接就能用。

变量	类型	含义
threadIdx	dim3	当前线程在块内的索引
blockIdx	dim3	当前线程块在网格内的索引
blockDim	dim3	每个线程块的维度（线程数）
gridDim	dim3	网格的维度（块数）

它们都是dim3类型，有.x、.y、.z三个分量。最常用的是.x，处理一维数据时够用了。

计算全局唯一ID

这是每个CUDA程序员必须掌握的公式：

int globalIdx = threadIdx.x + blockIdx.x * blockDim.x;

为什么是这个公式？你想想看：每个块有blockDim.x个线程，当前是第blockIdx.x个块，所以当前线程的全局索引就是块内索引 + 块编号 * 每块线程数。

处理二维数据时，公式会复杂一点：

int col = threadIdx.x + blockIdx.x * blockDim.x;
int row = threadIdx.y + blockIdx.y * blockDim.y;
int globalIdx = row * width + col;  // 行优先存储

我在做期权定价模型时，经常用二维网格来处理标的资产价格和到期时间两个维度。嗯，这样每个线程算一个价格点，并行效率很高。

边界检查

数据长度不一定能被线程数整除。所以每个核函数里都要做边界检查：

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

这个if (idx < N)看似简单，但少了它，程序可能崩溃，或者写出脏数据。我曾经在回测框架里漏了边界检查，结果回测结果偶尔对、偶尔错，排查了很久才发现是越界写坏了内存。

4.4 知识体系总览

下面这张图总结了本章的核心内容，帮你理清思路：

---

公众号:蓝海数据掘金营，微信deep3321