cuda-c-basics

CUDA C 编程基础

1. 核心概念

内核 (Kernel)

• 定义: 使用 __global__ 修饰的 C/C++ 函数，在 GPU 上并行执行
• 特点: 每个内核实例处理数据的一个子集，通过线程索引区分
• 调用: 使用 <<<grid_size, block_size>>> 语法从主机代码启动

网格 (Grid) 与块 (Block)

• 网格: 内核启动时的并行维度配置，如 (num_blocks_x, num_blocks_y)
• 块: 每个线程块包含的线程数，如 block_size = 256
• 关系: grid_size = ceil(total_elements / block_size)
• 限制: 每块最多 1024 线程（大多数 GPU）

线程层次

• Grid: 所有线程块的集合
• Block: 一组可以协作的线程（共享内存、同步）
• Warp: 32 个线程为一组并行执行（SIMT 执行模型）
• Thread: 最基本的执行单元

内存层次

• 全局内存 (Global Memory): 所有线程可访问，延迟高，容量大
• 共享内存 (Shared Memory): 块内线程共享，延迟低，容量有限（通常 48-164 KB/SM）
• 寄存器 (Registers): 每个线程私有，最快访问
• 常量内存 (Constant Memory): 只读，缓存优化
• 纹理内存 (Texture Memory): 只读，空间局部性优化

2. 标准内核结构（五步模式）

所有 CUDA C 内核都遵循相同的五步结构模式：

__global__ void standard_kernel(
    float* output, float* input, int n_elements
) {
    // 1. 计算全局线程索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 2. 边界检查
    if (idx < n_elements) {
        // 3. 加载数据
        float data = input[idx];
        
        // 4. 执行计算
        float result = compute_function(data);
        
        // 5. 存储结果
        output[idx] = result;
    }
}

内核启动方式

void launch_kernel(float* input, float* output, int n_elements) {
    const int block_size = 256;
    const int num_blocks = (n_elements + block_size - 1) / block_size;
    
    kernel<<<num_blocks, block_size>>>(output, input, n_elements);
}

3. 全局索引计算

一维数据处理

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

二维数据处理

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;

三维数据处理

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

网格配置

// 一维网格
int block_size = 256;
int num_blocks = (n_elements + block_size - 1) / block_size;
kernel<<<num_blocks, block_size>>>(...);

// 二维网格（矩阵操作）
dim3 block_size(16, 16);
dim3 grid_size((N + 15) / 16, (M + 15) / 16);
kernel<<<grid_size, block_size>>>(...);

// 三维网格（体积数据）
dim3 block_size(8, 8, 8);
dim3 grid_size((X + 7) / 8, (Y + 7) / 8, (Z + 7) / 8);
kernel<<<grid_size, block_size>>>(...);