Merge pull request 'optimize readme.md & add GPUCodeForces赛题解读.md' (#4) from Kuohais/GPUCodeForces:dev into main

This commit is contained in:
topshare 2025-09-23 13:32:51 +08:00
commit c521545f49
2 changed files with 244 additions and 84 deletions

View File

@ -0,0 +1,165 @@
# 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
该文件是赛题的 “基础参照系”,定义了任务边界、输入数据生成规则和标准 GTGround 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 激活函数的计算无需额外初始化参数(如权重、偏置),简化了模型初始化逻辑,聚焦核心的激活函数计算。
### (二)模块 2CUDA 优化实现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=256CUDA 硬件的经典配置1 个 Block 包含 256 个线程,可完美适配 GPU 的 Warp 调度(每个 Warp 含 32 个线程256=8×32num_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 环境下的报错;
数据与模型迁移:将输入张量和模型均移动到 GPUcuda(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
```
### (四)模块 4Prompt 设计参考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计算逻辑仍能覆盖所有元素。

163
README.md
View File

@ -23,115 +23,110 @@
## 📥 参赛流程
* 进入[GPUCodeForces赛事首页](https://gitlink.com/ccf-ai-infra/GPUCodeForces)登录参与本期比赛的Gitee账号完成一份成功合并到仓库内的提交即为参赛成功时间自由方法自由只要有灵感就可以动手开code~
一句话概括:进入[GPUCodeForces赛事首页](https://gitlink.com/ccf-ai-infra/GPUCodeForces)完成一份成功合并到仓库内的提交即为参赛成功时间自由方法自由只要有灵感就可以动手开code~
### 🌰举个栗子
* 登录or注册自己的Gitlink账号后进入赛事首页查看仓库内的文件内容。仔细阅读[how-to-contribute.md](https://gitlink.com/ccf-ai-infra/GPUCodeForces/tree/main/how-to-contribute.md)完成CLA签署并熟悉提交流程。
* 登录or注册自己的Gitlink账号后进入赛事首页初步查看仓库内的文件内容:
* 看到仓库内文件有一个example文件夹
.
├── S1(说明:第一季比赛名称,该目录下文件需选手自建)
│ ├── issue id 1(选手提交算子前创建的issue对应的id)
│ │ ├── cudacode.py(必要提交文件)
│ │ ├── torchcode.py(必要提交文件)
│ │ ├── run_code.py(必要提交文件)
│ │ ├── prompt.txt(必要提交文件)
│ ├── issue id 2
│ ├── issue id ...
├── example(样例提供,供大家上手)
│ ├── 001-example(文件结构与 issue id x 一致)
│ ├── 002-example
├── images(图片文件夹,参赛选手可忽略)
├── FAQ.md(社区收集的问题与解答)
├── LICENSE(证书,参赛选手可忽略)
├── README.md(赛题baseline)
├── how-to-contribute.md(提交指南看这里)
<img src="./images/readme_sample_check.png">
* 在S1文件夹的文件即为参赛选手需要提交的文件其余文件夹和文件皆为辅助参赛选手了解比赛、提供思路、排忧解难之用。
这是我们提供的一个样例,接下来我们在这个基础上进行一次完整的算子优化的提交(我们鼓励大家自己找到更好的算子并优化)。
* 我们将当前赛题的项目clone到本地电脑上建议使用git clone + 链接的方式在S1文件夹下创建一个以自己issue id命名的文件夹待提交的参赛代码文件都放在这里面
* 我们将样例clone到自己电脑上
并关注四份文件: torchcode.py、prompt.txt、cudacode_ori.py、example_cudacode.py最终需要提交的代码文件正是这四个。本次比赛在[模力方舟](https://ai.gitee.com/compute)平台上使用算力券购买容器实例:
<img src="./images/readme_git_compute.png">
接着便可以在云端实例上进行代码修改。相关算力券的领取方式请见[算力平台使用说明](https://ai.gitee.com/docs/compute/container)、[算力券兑换发放和兑换](https://ai.gitee.com/docs/billing/coupons)。
* 然后在该比赛仓库新建一个issue填写赛题。这里我们是对example-001算子优化因此issue的主题就可以是“对001-example数据集进行性能优化”
<img src="./images/readme_sample_issue.png">
可以看到这里有一个“#2”这是issue id你的算子优化、新算子都应该绑定一个独立的issue id最终有多少份issue被审核通过就表示提交成功了多少份。在即将提交的时候在该赛题仓库的S1文件夹下新建一个以该id命名无需带#号的文件夹该文件夹内容为四份必要文件和其他视参赛者情况需要补充的材料如readme文件、用到的其他数据集等
<img src="./images/readme_sample_folder.png">
* 准备工作就绪接下来看到example-001内的代码
```
Tipsissue id ≠ issue名称名称建议为对该算子的概括性描述。创建了issue后链接末尾处的数字即为issue id。同时也可以在issue界面的醒目位置查看如"#123"这样的数字标识。
```
* 做好了准备工作就可以开始尽情发挥去寻找算子亦或是优化算子。在example文件夹中提供了算子样例如果想简单上手可以查看该文件夹。虽然有效的算子优化也算一次提交但我们鼓励大家发现新的算子✨真正与其他选手拉开差距。
🔧简单介绍一下样例代码:
**example_torchcode.py** 基准模型Baseline。示例提供一个简单的PyTorch模型只包含一个ReLU激活函数。
* <span style="background-color: grey; color: black; user-select: none;">get_inputs()</span>:生成模型运行时需要的输入数据。
* <span style="background-color: grey; color: black; user-select: none;">get_init_inputs()</span>:成模型初始化所需的参数(这里就是权重矩阵 weight
<span style="background-color: grey; color: black; user-select: none;">example_cudacode.py</span>优化模型。示例使用PyTorch的load_inline功能直接编译和加载CUDA代码创建了一个新的模型类使用自定义CUDA实现替代PyTorch的ReLU。
**example_cudacode.py**优化模型。示例使用PyTorch的load_inline功能直接编译和加载CUDA代码创建了一个新的模型类使用自定义CUDA实现替代PyTorch的ReLU。
<span style="background-color: grey; color: black; user-select: none;">run_code.py</span>验证和性能测试脚本。验证自定义CUDA实现与原始PyTorch实现的数值精度一致性比较两种实现的性能计算加速比。
**run_code.py**验证和性能测试脚本。验证自定义CUDA实现与原始PyTorch实现的数值精度一致性比较两种实现的性能计算加速比。
<span style="background-color: grey; color: black; user-select: none;">prompt.txt</span>:这里给予参赛者一些提示:
**prompt.txt**:提示词文本。提供类似 “融合算子 CUDA 设计” 的 prompt 编写思路。
* 要求编写自定义CUDA内核来替换PyTorch算子以获得加速
💡如何将一个example变为自己的一份提交具体的算子优化思路可参考[GPUCodeForces赛题解读](./GPUCodeForces赛题解读.md)
* 为了测试代码最终的跑通结果需要使用规定的GPU。在[模力方舟](https://ai.gitee.com/compute)平台上准备了大家此次需要的算力资源,使用免费的算力券购买实例,接着便可以在云端实例上进行代码修改和测试。
相关算力券的领取方式请见[算力平台使用说明](https://ai.gitee.com/docs/compute/container)、[算力券兑换发放和兑换](https://ai.gitee.com/docs/billing/coupons)。
```
Tips进行云端服务器的实例选择有库存的即可点击进入后需要确定配置。其余选项保持默认只需更改镜像处的选择基础镜像-->CV-CUDA-->PyTorch2.4.0-->Python任意-->maca 3.0.0.5
```
* 可以自由选择替换哪些算子,考虑算子融合机会
* 代码测试运行没有问题后便可以准备提交了记得将所有必要文件保存到提交文件夹。提交流程如下全程使用git
```
# cd到自己克隆到本地的项目路径下C:\Users\Desktop\ODTC AI Infra\GPUCodeForces
* 提供了示例语法和内联嵌入自定义CUDA算子的方法
* 给出了需要优化的模型架构简单的ReLU模型
git remote -v # 检查是否链接到自己的Gitlink仓库
git checkout -b dev # 新建一个分支dev即分支名字可以任取
git add . # 将文件的变更操作暂存
git commit -m "输入你本次提交的目的" # 目的简洁明了最好
git push origin dev # 提交你的修改到分支上
```
对git与远程仓库操作的疑问可以点击[how-to-contribute.md](https://gitlink.org.cn/ccf-ai-infra/GPUCodeForces/tree/main/how-to-contribute.md)查看详情。
* 然后参照example-001文件夹创建自己的文件夹提出新的torch cude对其中torch可来源于PyTorch、PaddlePaddle、TensorFlow、Jax、MMCV、Transformers等框架可以是单个算子也可以是多个算子融合。cuda代码可自己编写或者参照prompt.txt让LLM辅助编写
1. 修改torchcode.py按照example_torchcode.py格式在Model的__init__和forward中放入目标torch代码注意保留get_inputs()和get_init_inputs()函数
2. 修改cudacode.py修改cudacode.py里的load_inline和ModelNewcuda代码核心实现放在source字段中
* 顺利提交后的代码还只在你自己fork的仓库下还需要和主仓库合并才能真正让管理员看到你的代码。
* 优化好后,可以在模力方舟的实例上运行:
fork仓库你自己的仓库与主仓库之间关系如下
```sh
python run_code.py
```
确保能够正确输出结果后再准备提交。
```
ODTC AI Infra/GPUCodeForces(main)
├── ...
├── folders
* 接下来将优化好的代码保存到本地,然后参照[how-to-contribute.md](https://gitlink.org.cn/ccf-ai-infra/GPUCodeForces/tree/main/how-to-contribute.md)的指引进行代码仓库的提交与合并。
your_name/GPUCodeForces(main) # 自己仓库下的main分支内容与主仓库一致
├── ...
├── folders
your_name/GPUCodeForces(dev) # dev分支才是你修改后的代码存放处
├── ...
│ ├── your codes
├── folders
```
进入自己的仓库,点击上方选项栏:
```
合并请求(PR)-->+新建合并请求-->源分支选择dev(名称任取)-->填写下方选框-->标题为自己的issue标题-->描述填请简明扼要描述内关联自己的issue id(如“fixes #001”)
```
提交后便能看到自己的PR记录了在对应记录的评论区会有测试结果的告知请留意查看~
* 最终,成功提交的代码会合并到 S1/#your_issue id 下并且你的相关pr也会关闭。就像下面这样
<img src="./images/readme_sample_merge.png">
🌳一份完整的提交流程如上,期待各位自由发挥,赛出风采与水平!
### 📦 提交PR内容
* **一个PR包含样本的目录** [提交样例](https://gitlink.org.cn/ccf-ai-infra/GPUCodeForces/tree/main/example/001-example)
* 每个提交目录建议包含如下:
1. **示例代码:** torch代码示例
2. **对比代码:** 和torch对应的CUDA代码
3. **测试代码入口:** run\_code.py请务必用这个名称提交的PR会根据这个名称在GPU上测试结果
4. **其它文件(或目录):** prompt利用LLM从torch代码生成cuda代码的prompt示例或者其它优化代码
5. **PR目录说明文件** https://gitlink.org.cn/ccf-ai-infra/GPUCodeForces/tree/main/example/001-example/readme.md
### 📦 提交PR的格式
建议在开始做题目之前创建一个赛题提交的PR和自己创建的赛题相关联。参赛选手在每个比赛周期的目录下例如第一期S1、第二期S2、第三期S3...创建一个目录目录名称赛题的IDICTXSZ),例如:
```plaintext
.
├── S1(说明:第一季比赛名称)
│ ├── ICTXSZ(说明以赛题ID命名的目录存放PR提交样本的目录)
| | ├── 示例代码
│ | ├── 对比代码
| | └── ……
│ └── ……
└── S2(第二季比赛)
└── 赛题1
```
⏺如仍有疑问,请点击[提交流程演示视频](https://www.bilibili.com/video/BV1CcnTztEfc/)
### ⭐审核流程
* 在一切文件都准备好并且提交后在对应的PR下会得到回复
<image src="./images/readme_comment_check.png">
也就是说,除了能够自己在服务器上运行得到算子测算的初步结果外,还可以在这里看到最终的测算结果。这里显示测试通过才能进入后续审核流程。
* 你提交的PR都会得到回复大概存在的几种情况如下
```
提交PR-->测试通过✔️-->已合并-->有效提交 ヽ(✿゚▽゚)
提交PR-->测试通过✔️-->已关闭-->代码重复/相似-->无效提交 (;′⌒`)
提交PR-->测试失败✖️-->已关闭-->代码不合格-->无效提交 (;′⌒`)
```
也就是说,除了能够自己在算力平台的实例上运行得到算子测算的初步结果外,还可以在这里看到最终的测算结果。这里显示测试通过才能进入后续审核流程,并最终上传至比赛的算子仓库。
### ✅ 参赛资格