cuda-c-patterns

CUDA C 编程模式

1. 向量操作模式

适用于元素级运算：加法、乘法、激活函数等。

标准代码结构

__global__ void vector_add_kernel(
    const float* a, const float* b, float* c, int n_elements
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < n_elements) {
        c[idx] = a[idx] + b[idx];
    }
}

适用算子

• 算术运算: add, mul, sub, div
• 激活函数: relu, sigmoid, tanh, gelu, silu
• 数学函数: exp, log, sqrt, pow, abs
• 类型转换: cast
• 广播操作: broadcast

关键要点

• 使用一维索引 blockIdx.x * blockDim.x + threadIdx.x
• 边界检查 if (idx < n_elements)
• 简单直接的数据流：加载 → 计算 → 存储
• 推荐块大小: 256 或 512

ReLU 示例

__global__ void relu_kernel(
    const float* input, float* output, int n
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = fmaxf(input[idx], 0.0f);
    }
}

GELU 示例

__global__ void gelu_kernel(
    const float* input, float* output, int n
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float x = input[idx];
        // 近似 GELU: 0.5 * x * (1 + tanh(sqrt(2/pi) * (x + 0.044715 * x^3)))
        float cdf = 0.5f * (1.0f + tanhf(0.7978845608f * (x + 0.044715f * x * x * x)));
        output[idx] = x * cdf;
    }
}

多输入逐元素操作

__global__ void fused_multiply_add_kernel(
    const float* a, const float* b, const float* c,
    float* output, int n
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        output[idx] = a[idx] * b[idx] + c[idx];
    }
}

2. 归约模式

适用于求和、最大值、最小值等聚合操作。

标准代码结构（共享内存归约）

__global__ void reduction_sum_kernel(
    const float* input, float* output, int n_elements
) {
    extern __shared__ float sdata[];
    
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 加载数据到共享内存
    sdata[tid] = (idx < n_elements) ? input[idx] : 0.0f;
    __syncthreads();
    
    // 块内归约（树形归约）
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            sdata[tid] += sdata[tid + s];
        }
        __syncthreads();
    }
    
    // 第一个线程写入块级结果
    if (tid == 0) {
        atomicAdd(output, sdata[0]);
    }
}