wenyi-li 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>>>(...);
4. 边界处理
基本边界检查
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n_elements) {
// 安全访问 input[idx]
output[idx] = input[idx];
}
二维边界检查
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < M && col < N) {
// 安全访问矩阵元素
output[row * N + col] = input[row * N + col];
}
5. 内存管理模式
主机-设备数据传输
// 分配设备内存
float* d_input, *d_output;
cudaMalloc(&d_input, size * sizeof(float));
cudaMalloc(&d_output, size * sizeof(float));
// 拷贝数据到设备
cudaMemcpy(d_input, h_input, size * sizeof(float), cudaMemcpyHostToDevice);
// 启动内核
kernel<<<grid, block>>>(d_output, d_input, size);
// 拷贝结果回主机
cudaMemcpy(h_output, d_output, size * sizeof(float), cudaMemcpyDeviceToHost);
// 释放内存
cudaFree(d_input);
cudaFree(d_output);
6. PyTorch 集成方式
⚠️ 重要：使用 Python JIT 编译方式
在生成算子时，必须在 Python 模块中内嵌 CUDA C 代码，使用 torch.utils.cpp_extension.load_inline 进行 JIT 编译：
import torch
from torch.utils.cpp_extension import load_inline
source = """
#include <torch/extension.h>
#include <cuda_runtime.h>
__global__ void my_kernel(const float* input, float* output, int size) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < size) {
output[idx] = input[idx] * 2.0f;
}
}
torch::Tensor my_kernel_call(torch::Tensor input) {
auto size = input.numel();
auto output = torch::zeros_like(input);
int block_size = 256;
int num_blocks = (size + block_size - 1) / block_size;
my_kernel<<<num_blocks, block_size>>>(
input.data_ptr<float>(), output.data_ptr<float>(), size);
return output;
}
"""
cpp_src = "torch::Tensor my_kernel_call(torch::Tensor input);"
kernel_module = load_inline(
name="my_cuda",
cpp_sources=cpp_src,
cuda_sources=source,
functions=["my_kernel_call"],
verbose=True,
extra_cflags=[""],
extra_ldflags=[""],
)
7. 最佳实践总结
编程原则
- 单一职责: 每个内核只做一件事
- 参数简单: 避免复杂的数据结构传递
- 边界检查: 所有数组访问前必须检查边界
- 内存局部性: 尽量访问相邻内存位置
常见错误避免
- 越界访问: 忘记边界检查导致运行时错误或结果异常
- 内存泄漏: 忘记调用
cudaFree释放设备内存 - 同步错误: 共享内存操作缺少
__syncthreads() - 类型不匹配: 主机和设备端数据类型不一致
- 设备内存不足: 块大小过大或数据量超出 GPU 内存
⚠️ 注意事项
- 生成的内核代码不要包含任何测试代码片段
- 禁止使用
printf()、throw std::runtime_error()等打印/异常语句 - 内核内不使用
malloc/new进行动态分配