diff --git a/GPUCodeForces赛题解读.md b/GPUCodeForces赛题解读.md new file mode 100644 index 0000000..6afabd0 --- /dev/null +++ b/GPUCodeForces赛题解读.md @@ -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） +该文件是赛题的 “基础参照系”，定义了任务边界、输入数据生成规则和标准 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.代码间逻辑流程图 + + + +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计算逻辑仍能覆盖所有元素。 diff --git a/README.md b/README.md index 7f425ec..e975374 100644 --- a/README.md +++ b/README.md @@ -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(提交指南看这里) - + * 在S1文件夹的文件即为参赛选手需要提交的文件，其余文件夹和文件皆为辅助参赛选手了解比赛、提供思路、排忧解难之用。 - 这是我们提供的一个样例，接下来我们在这个基础上进行一次完整的算子优化的提交（我们鼓励大家自己找到更好的算子并优化）。 + * 我们将当前赛题的项目clone到本地电脑上（建议使用git clone + 链接的方式），在S1文件夹下创建一个以自己issue id命名的文件夹，待提交的参赛代码文件都放在这里面。 - * 我们将样例clone到自己电脑上， - 并关注四份文件： torchcode.py、prompt.txt、cudacode_ori.py、example_cudacode.py，最终需要提交的代码文件正是这四个。本次比赛在[模力方舟](https://ai.gitee.com/compute)平台上使用算力券购买容器实例： - - - - 接着便可以在云端实例上进行代码修改。相关算力券的领取方式请见[算力平台使用说明](https://ai.gitee.com/docs/compute/container)、[算力券兑换发放和兑换](https://ai.gitee.com/docs/billing/coupons)。 - - * 然后在该比赛仓库新建一个issue，填写赛题。这里我们是对example-001算子优化，因此issue的主题就可以是“对001-example数据集进行性能优化”： - - - - 可以看到这里有一个“#2”，这是issue id，你的算子优化、新算子都应该绑定一个独立的issue id（最终有多少份issue被审核通过，就表示提交成功了多少份）。在即将提交的时候，在该赛题仓库的S1文件夹下新建一个以该id命名（无需带#号）的文件夹，该文件夹内容为四份必要文件和其他视参赛者情况需要补充的材料（如readme文件、用到的其他数据集等）： - - - - * 准备工作就绪，接下来看到example-001内的代码： + ``` + Tips：issue id ≠ issue名称，名称建议为对该算子的概括性描述。创建了issue后，链接末尾处的数字即为issue id。同时，也可以在issue界面的醒目位置查看如"#123"这样的数字标识。 + ``` + * 做好了准备工作，就可以开始尽情发挥，去寻找算子亦或是优化算子。在example文件夹中提供了算子样例，如果想简单上手可以查看该文件夹。虽然有效的算子优化也算一次提交，但我们鼓励大家发现新的算子✨真正与其他选手拉开差距。 + + 🔧简单介绍一下样例代码： **example_torchcode.py：** 基准模型（Baseline）。示例提供一个简单的PyTorch模型，只包含一个ReLU激活函数。 - - * get_inputs()：生成模型运行时需要的输入数据。 - * get_init_inputs()：成模型初始化所需的参数（这里就是权重矩阵 weight）。 - example_cudacode.py：优化模型。示例使用PyTorch的load_inline功能直接编译和加载CUDA代码，创建了一个新的模型类，使用自定义CUDA实现替代PyTorch的ReLU。 + **example_cudacode.py**：优化模型。示例使用PyTorch的load_inline功能直接编译和加载CUDA代码，创建了一个新的模型类，使用自定义CUDA实现替代PyTorch的ReLU。 - run_code.py：验证和性能测试脚本。验证自定义CUDA实现与原始PyTorch实现的数值精度一致性，比较两种实现的性能，计算加速比。 + **run_code.py**：验证和性能测试脚本。验证自定义CUDA实现与原始PyTorch实现的数值精度一致性，比较两种实现的性能，计算加速比。 - prompt.txt：这里给予参赛者一些提示： + **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和ModelNew，cuda代码核心实现放在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也会关闭。就像下面这样： - - 🌳一份完整的提交流程如上，期待各位自由发挥，赛出风采与水平！ - -### 📦 提交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...）创建一个目录，目录名称赛题的ID（ICTXSZ)，例如： - -```plaintext -. -├── S1(说明：第一季比赛名称) -│ ├── ICTXSZ(说明：以赛题ID命名的目录存放PR提交样本的目录) -| | ├── 示例代码 -│ | ├── 对比代码 -| | └── …… -│ └── …… -└── S2(第二季比赛) - └── 赛题1 - -``` + ⏺如仍有疑问，请点击[提交流程演示视频](https://www.bilibili.com/video/BV1CcnTztEfc/) ### ⭐审核流程 -* 在一切文件都准备好并且提交后，在对应的PR下会得到回复： - - - - -也就是说，除了能够自己在服务器上运行得到算子测算的初步结果外，还可以在这里看到最终的测算结果。这里显示测试通过才能进入后续审核流程。 +* 你提交的PR都会得到回复，大概存在的几种情况如下 + ``` + 提交PR-->测试通过✔️-->已合并-->有效提交 ヽ(✿ﾟ▽ﾟ)ノ + 提交PR-->测试通过✔️-->已关闭-->代码重复/相似-->无效提交 (；′⌒`) + 提交PR-->测试失败✖️-->已关闭-->代码不合格-->无效提交 (；′⌒`) + ``` + 也就是说，除了能够自己在算力平台的实例上运行得到算子测算的初步结果外，还可以在这里看到最终的测算结果。这里显示测试通过才能进入后续审核流程，并最终上传至比赛的算子仓库。 ### ✅ 参赛资格