forked from ccf-ai-infra/GPUCodeForces
166 lines
11 KiB
Markdown
166 lines
11 KiB
Markdown
# GPUCodeForces赛题解读
|
||
## 一、赛题核心定位与整体框架
|
||
本赛题属于GPU CUDA 性能优化类任务，要求参赛选手通过自定义 CUDA Kernel 实现各类函数，并与 PyTorch 内置实现进行精度对齐和性能比拼。
|
||
这份赛题解读，我们以“通过自定义CUDA Kernel实现Swish激活函数”来作为引子，让大家从一个具体例子中了解到算子优化的细节。
|
||
整套关联代码（example_torchcode.py、example_cudacode.py、run_code.py）构成了 “任务定义 - 基准实现 - CUDA 优化 - 评测验证” 的完整闭环，prompt.txt则提供了类似 “融合算子 CUDA 设计” 的 prompt 编写思路，可作为加分项参考。其核心目标是考察选手的 CUDA 内核设计能力、内存效率优化能力及精度与性能的平衡能力。
|
||
|
||
## 二、赛题模块拆解与代码解读
|
||
### （一）模块 1：任务定义与基准实现（example_torchcode.py）
|
||
该文件是赛题的 “基础参照系”，定义了任务边界、输入数据生成规则和标准 GT（Ground Truth）输出，对应评测数据集要求中的 “数据集样本描述”“输入数据生成函数”“标准 GT 输出生函数”。
|
||
1. 任务定义：Swish 激活函数计算
|
||
|
||
Swish 是深度学习中优于 ReLU 的激活函数，数学表达式为：Swish(x) = x * sigmoid(x)。其中sigmoid(x) = 1 / (1 + exp(-x))，核心作用是为模型引入非线性，且在大维度张量（如隐藏层特征）上的计算效率直接影响模型整体推理速度。代码中通过Model类实现基准逻辑：
|
||
|
||
```
|
||
class Model(nn.Module):
|
||
def forward(self, x: torch.Tensor) -> torch.Tensor:
|
||
return x * torch.sigmoid(x) # 标准GT实现
|
||
```
|
||
|
||
2. 输入数据生成函数
|
||
|
||
赛题定义了固定的输入规模（模拟大模型隐藏层维度，具有实际业务代表性），函数get_inputs()生成符合任务需求的输入张量：
|
||
|
||
```
|
||
batch_size = 16 # 批量大小
|
||
dim = 16384 # 单样本特征维度（模拟大模型隐藏层）
|
||
def get_inputs():
|
||
x = torch.randn(batch_size, dim) # 随机正态分布张量（符合深度学习输入特性）
|
||
return [x]
|
||
```
|
||
|
||
输入特性：torch.randn生成均值为 0、方差为 1 的浮点数张量，覆盖正负值，能全面验证 Swish 在不同输入下的计算精度；
|
||
|
||
规模选择：16×16384的张量共 262,144 个元素，既能暴露 CUDA 线程调度的效率差异，又不会因规模过大导致测试耗时过长。
|
||
|
||
3. 初始化输入函数
|
||
|
||
get_init_inputs()返回空列表，因 Swish 激活函数的计算无需额外初始化参数（如权重、偏置），简化了模型初始化逻辑，聚焦核心的激活函数计算。
|
||
|
||
### （二）模块 2：CUDA 优化实现（example_cudacode.py）
|
||
|
||
该文件是参赛选手的 “核心提交目标”，对应评测数据集要求中的 “CUDA 解决方案实现”，需通过自定义 CUDA Kernel 实现 Swish 函数，同时满足精度和性能要求。
|
||
|
||
1. 实现思路：轻量级高效 Kernel 设计
|
||
|
||
Swish 是逐元素操作（每个输出元素仅依赖对应输入元素），无需跨元素通信，因此 Kernel 设计聚焦 “线程调度效率” 和 “浮点计算稳定性”：
|
||
|
||
线程索引：采用 1D 索引（idx = blockIdx.x * blockDim.x + threadIdx.x），每个线程处理 1 个元素，避免复杂的 2D/3D 索引计算开销；
|
||
|
||
Block/Grid 配置：block_size=256（CUDA 硬件的经典配置，1 个 Block 包含 256 个线程，可完美适配 GPU 的 Warp 调度（每个 Warp 含 32 个线程，256=8×32）），num_blocks通过 “向上取整” 计算（(size + block_size - 1) / block_size），确保覆盖所有元素；
|
||
|
||
浮点计算：使用expf()（单精度浮点数指数函数），与输入张量的float32类型匹配，避免精度浪费和类型转换开销。
|
||
|
||
2. 代码编译与封装
|
||
|
||
通过torch.utils.cpp_extension.load_inline实现内联 CUDA 代码编译，无需单独编写setup.py，简化开发流程：
|
||
|
||
```
|
||
swish = load_inline(
|
||
name="swish", # 模块名
|
||
cpp_sources=swish_cpp_source, # C++声明（接口定义）
|
||
cuda_sources=swish_source, # CUDA内核实现
|
||
functions=["swish_cuda"], # 暴露给Python的函数
|
||
verbose=True # 打印编译日志（便于调试）)
|
||
```
|
||
|
||
并通过ModelNew类封装 CUDA 函数，保持与example_torchcode.py中Model类一致的接口（forward方法），确保后续评测代码可无缝调用。
|
||
|
||
### （三）模块 3：精度与性能评测（run_code.py）
|
||
该文件是赛题的 “评测执行器”，对应评测数据集要求中的 “性能评估指标” 和 “正确性验证”，实现了从数据准备到结果分析的全流程自动化评测。
|
||
|
||
1. 评测前置准备
|
||
|
||
CUDA 可用性检查：先判断torch.cuda.is_available()，避免无 GPU 环境下的报错；
|
||
|
||
数据与模型迁移：将输入张量和模型均移动到 GPU（cuda(device=device)），确保计算在 GPU 上执行；
|
||
|
||
GPU 预热：先执行 10 次空计算（_ = torch_model(*inputs)），避免 GPU 初始化、内存分配等一次性开销影响性能计时精度。
|
||
|
||
2. 精度对齐验证（核心指标）
|
||
|
||
精度是 CUDA 实现的 “准入条件”，需确保自定义 Kernel 与 PyTorch 基准的输出误差在可接受范围：
|
||
```
|
||
abs_diff = torch.abs(output_torch - output_cuda) # 逐元素绝对误差
|
||
max_diff = torch.max(abs_diff).item() # 最大误差（全局）
|
||
mean_diff = torch.mean(abs_diff).item() # 平均误差（全局）
|
||
```
|
||
|
||
误差阈值：max_diff < 1e-4且mean_diff < 1e-5，符合深度学习中浮点计算的常见容忍度（单精度浮点数的机器 epsilon 约为 1e-7，该阈值留有充足余量）；
|
||
|
||
结果判定：若满足阈值则 “精度对齐”，否则视为无效实现，无法进入性能评测环节。
|
||
|
||
3. 性能加速比测试（核心指标）
|
||
|
||
性能评测聚焦 “平均执行时间” 和 “加速比”，对应评测数据集要求的 “执行时间”“吞吐量” 指标：迭代次数：100 次（减少随机波动，确保计时稳定性）；同步计时：使用torch.cuda.synchronize()强制等待 GPU 计算完成，避免 CUDA 异步执行导致的计时偏差；结果计算：平均时间 = 总时间 / 迭代次数（消除单次执行的偶然误差）；加速比 = PyTorch 平均时间 / CUDA 平均时间（比值越大，CUDA 优化效果越好）。
|
||
|
||
4. 评测输出示例
|
||
```
|
||
-------------------- 精度对齐验证 --------------------
|
||
✅ 精度对齐：最大误差 0.000089，平均误差 0.000012
|
||
-------------------- 性能加速比测试 --------------------
|
||
PyTorch内置Swish平均执行时间: 0.000123秒
|
||
自定义CUDA Swish平均执行时间: 0.000045秒
|
||
加速比 (Speedup): 2.73x
|
||
```
|
||
### （四）模块 4：Prompt 设计参考（prompt.txt）
|
||
|
||
针对 “矩阵乘法 + GELU 融合算子”，提供了赛题 “加分项 Prompt” 的设计思路，对应评测数据集要求的 “加分项：LLM 生成 CUDA 代码的 Prompt”。
|
||
|
||
1. Prompt 设计核心要素
|
||
|
||
任务拆解：明确 “先矩阵乘、后 GELU” 的原始流程，指出融合的必要性（避免中间结果的全局内存读写）；
|
||
|
||
技术要求：指定关键优化点（2D Grid/Block、共享内存 tiling），引导 LLM 生成符合 CUDA 最佳实践的代码；
|
||
|
||
精度约束：强调 “数值稳定性和精度”，避免 LLM 为追求性能牺牲精度。
|
||
|
||
2. 迁移应用到 Swish 任务
|
||
|
||
若为 Swish 设计 Prompt，可参考如下结构：
|
||
```
|
||
Write a custom CUDA kernel for Swish activation (Swish(x) = x * sigmoid(x)).
|
||
The original PyTorch implementation uses x * torch.sigmoid(x), which may have redundant global memory access.
|
||
You should optimize the CUDA kernel to:
|
||
- Use 1D grid/block dimensions (each thread processes one element)
|
||
- Choose appropriate block size (e.g., 256) for GPU warp scheduling
|
||
- Ensure numerical stability (use float32 and expf() for sigmoid)
|
||
The input is a PyTorch tensor of shape (batch_size=16, dim=16384), and the output should match PyTorch's result with max error < 1e-4.
|
||
```
|
||
|
||
## 三、赛题核心考察点
|
||
|
||
CUDA Kernel 设计能力：线程索引计算、Block/Grid 配置合理性（如block_size=256的选择依据）；
|
||
|
||
精度控制能力：浮点计算稳定性（如匹配float32类型、避免exp()溢出）；
|
||
|
||
性能优化意识：GPU 预热、同步计时、减少冗余内存访问；
|
||
|
||
工程化实现能力：CUDA 代码与 PyTorch 的接口兼容（如ModelNew类的封装）、编译调试能力。
|
||
|
||
## 四、代码间逻辑流与参赛指引
|
||
|
||
1.代码间逻辑流程图
|
||
|
||
<img src="./images/code_step.png">
|
||
|
||
2.参赛选手操作指引
|
||
|
||
参考example_torchcode.py理解任务边界（输入规模、GT 输出）；
|
||
|
||
编写自定义 CUDA Kernel（可借鉴example_cudacode.py的结构，优化 Block/Grid 或浮点计算）；
|
||
|
||
使用run_code.py验证精度（需满足误差阈值）和性能（追求更高加速比）；
|
||
|
||
（加分项）设计 Prompt，让 LLM 生成你的 CUDA 代码，并对比 LLM 生成结果与手写结果的差异。
|
||
|
||
## 五、关键注意事项
|
||
|
||
精度优先：若 CUDA 实现性能极高但精度不达标，视为无效提交；
|
||
|
||
数据类型统一：输入张量、CUDA 计算均使用float32，避免float64导致的性能下降；
|
||
|
||
计时准确性：必须使用torch.cuda.synchronize()，否则异步执行会导致计时结果偏小（虚假性能提升）；
|
||
|
||
可扩展性：若输入规模变化（如dim=32768），需确保num_blocks计算逻辑仍能覆盖所有元素。
|