diff --git a/README.md b/README.md index b0360381c4f69a5ef41f2ee8d658b994e5e8bcf6..096853fab255cd1c8b59540786ea49825e144106 100644 --- a/README.md +++ b/README.md @@ -1,80 +1,204 @@ # 评测数据集生成挑战赛（GPUCodeForces） -欢迎参加 **评测数据集生成挑战赛** 📊！ -本比赛旨在构建一个标准化、可用于 GPU 性能测试的评测数据集，帮助开发者更高效地比较不同硬件和框架的性能表现。 +欢迎参加 **评测数据集生成挑战赛** 📊！   本比赛旨在构建一个标准化、可用于 GPU 性能测试的评测数据集，帮助开发者更高效地比较不同硬件和框架的性能表现。 --- -## 🧠 比赛背景简介 +## 🧠 比赛背景简介 -在 AI 模型开发和部署中，**GPU 性能评测**是一个非常重要的环节。 -不同 GPU、不同深度学习框架（如 PyTorch、TensorFlow、PaddlePaddle 等）在运行相同任务时，速度、吞吐量、内存占用等表现差异很大。 -本次挑战赛希望通过社区的力量，**构建一个标准化、带权重的评测数据集**，让 GPU 性能比较更加科学、公正。 +在 AI 模型开发和部署中，**GPU 性能评测**是一个非常重要的环节。   不同 GPU、不同深度学习框架（如 PyTorch、TensorFlow、PaddlePaddle 等）在运行相同任务时，速度、吞吐量、内存占用等表现差异很大。   本次挑战赛希望通过社区的力量，**构建一个标准化、带权重的评测数据集**，让 GPU 性能比较更加科学、公正。 --- -## 🎯 比赛目标 +## 🎯 比赛目标 -- 从 **PyTorch、PaddlePaddle、TensorFlow、Jax、MMCV、Transformers** 等框架中收集并生成评测样本。 -- 为每个样本提供**标准输出**和**性能指标**，确保结果可复现。 -- 最终形成 **GPU CodeForces** 数据集和评价方法。 +* 从 **PyTorch、PaddlePaddle、TensorFlow、Jax、MMCV、Transformers** 等框架中收集并生成评测样本。 + +* 为每个样本提供**标准输出**和**性能指标**，确保结果可复现。 + +* 最终形成 **GPU CodeForces** 数据集和评价方法。 + --- -## 📥 如何参与提交？ +## 📥 如何参与提交？ -### ✅ 参赛资格 -- 你提交的样本中，**至少有 1 个样本被评审通过并正式整合到“GPU CodeForces” 数据集**，即可算作有效参赛。 +### 📝初始化仓库及提交流程 -### 📦 提交内容 -- **一个 JSON 格式的数据集文件**（详细格式请参考《评测数据集生成挑战赛样本和要求说明》）。 -- 每个数据集需包含： - 1. **测试样本**（代码或数据） - 2. **标准 GT（Ground Truth）输出生成函数**（基于 CPU/Numpy 或原始框架） - 3. **CUDA 性能评估指标**（执行时间、吞吐量、内存带宽） +* 进入到赛事首页：[GPUCodeForces赛事首页](https://gitee.com/ccf-ai-infra/GPUCodeForces)，点击[Issues](https://gitee.com/ccf-ai-infra/GPUCodeForces/issues/new/choose)，选择赛题\[CP\]->立即开始： + ---- +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/1.png) -## 📈 评分规则 +然后按照模板格式填写相关信息后提交即可。在最后一个选项“推荐其他选手完成”这里，若有团队则可框选，若是自身完成则不用框选： -### 📊 基础得分 -| 内容 | 分值 | -|------|------| -| 提供标准 GT 输出生成函数（Numpy-CPU / 原始框架实现） | +2 分 | -| CUDA 执行时间评估 | +5 分 | -| CUDA 吞吐量评估 | +4 分 | -| CUDA 内存带宽评估 | +3 分 | +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/2.png) -### ✨ 加分项 -| 内容 | 分值 | -|------|------| -| 提供 Prompt 让 LLM 生成对应的 CUDA 代码，并同样进行性能评价 | 额外加分 | +提交完成后，返回Issue主页可以看到自己的提交记录，可以复制Issue id进行保存： + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/3.png) + +* 回到赛事主页，在仓库界面右上角选择fork至自己的仓库下： + + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/4.png) + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/5.png) + +然后把源项目克隆至本地，文件路径可以根据自己的喜好来： + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/6.png) + +在克隆仓库或本地文件夹内都可以大展你的身手，进行算子优化了🎉！ + +* 准备好了代码及各项必需文件后，在本地文件夹内启动git bash，进行初始化和远程仓库连接： + + +```shell +git init +git remote add origin https://gitee.com/your-name/GPUCodeForces.git +git remote -v +#成功后应显示： +#origin https://gitee.com/your-name/GPUCodeForces.git (fetch) +#origin https://gitee.com/your-name/GPUCodeForces.git (push) +``` + +* 然后将本地优化好的算子仓库存入到暂存区，并进行上传： + + +```shell +git add . +git commit -m "填写本次提交记录，如'第一次提交'" +git push -u origin master +#最后一步需要根据自身的实际仓库分支来，一般为master，也可能为main或者你自定义的目录名称，可以使用 git push --upstream orgin 目录名 +``` + +* 上传成功后，回到自己的仓库主页，选择”新建Pull Request“： + + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/7.png) + +然后在上传界面选择自己更新项目的分支，选择合并即可： + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/8.png) + +### 📦 提交PR内容 + +* **一个PR包含样本的目录** [提交样例](https://gitee.com/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://gitee.com/ccf-ai-infra/GPUCodeForces/blob/main/example/001-example/readme.md + + +### 📦 提交PR的格式 + +建议在开始做题目之前创建一个赛题，提交的PR和自己创建的赛题相关联。参赛选手在每个比赛周期的目录下（例如：第一期S1、第二期S2、第三期S3...）创建一个目录，目录名称赛题的ID（ICTXSZ)，例如： -> **接受数量** = 提交并被评审通过的样本总数 +```plaintext +. +├── S1(说明：第一季比赛名称) +│ ├── ICTXSZ(说明：以赛题ID命名的目录存放PR提交样本的目录) +| | ├── 示例代码 +│ | ├── 对比代码 +| | └── …… +│ └── …… +└── S2(第二季比赛) + └── 赛题1 + +``` + +### 💻审核流程 + +* 在一切文件都准备好并且提交后，在对应的PR下会得到回复： + + +![image.png](https://gitee.com/broad-sea-life/test-for-template/raw/master/Images/9.png) + +也就是说，除了能够自己在服务器上运行得到算子测算的初步结果外，还可以在这里看到最终的测算结果。这里显示测试通过才能进入后续审核流程。 + +### ✅ 参赛资格 + +* 你提交的PR样本中，**至少有 1 个样本被评审通过并正式整合到“GPU CodeForces” 数据集**，即可算作有效参赛。 + + +### ⚠️注意事项 + +1. 请勿抄袭他人代码或成果 + +2. 请勿恶意提交（如相同算子多次提交、相近优化重复提交） + \---相近优化：即指同一份参赛用例在优化后加速比没有提升或更低的用例提交 + +3. 请遵守提交的格式要求、内容规范 + --- -## 🏅 排名机制 +## 🏅 竞赛排名机制 + +1. **优先按接受数量从高到低排序作为排名，取前12名。** + +2. 若接受数量相同： + + * 比较总基础评分高者优先 + + * 若仍相同，比加分项得分高者优先 + + +> **接受数量** = 提交并被评审通过的样本总数 -1. **优先按接受数量从高到低排序** -2. 若数量相同： - - 比较总评分高者优先 - - 若仍相同，比加分项得分高者优先 +> **接受数量相同需要区分排名时如下的基础和甲方的评分规则才会生效** --- -## 📚 术语解释 +## 📈 评分规则 -- **评测数据集**：用来测试 GPU 性能的一组标准化样本，包括代码、输入数据和预期结果。 -- **GT（Ground Truth）**：标准参考答案或结果，用来验证程序运行是否正确。 -- **吞吐量（Throughput）**：每秒钟能处理的数据量，越高表示 GPU 处理能力越强。 -- **内存带宽（Memory Bandwidth）**：单位时间内 GPU 内存与计算核心之间的数据传输速度。 -- **Prompt**：引导大语言模型（LLM）生成代码或内容的提示词。 -- **LLM**：Large Language Model，大语言模型，如 ChatGPT、LLaMA 等。 +### 📊 基础得分 + +| 内容 | 分值 | +| --- | --- | +| 提供标准 GT 输出生成函数（Numpy-CPU / 原始框架实现） | +2 分 | +| CUDA 执行时间评估 | +5 分 | +| CUDA 吞吐量评估 | +4 分 | +| CUDA 内存带宽评估 | +3 分 | + +### ✨ 加分项 + +| 内容 | 分值 | +| --- | --- | +| 提供 Prompt 让 LLM 生成对应的 CUDA 代码，并同样进行性能评价 | 额外加分 | --- -## 📬 联系与帮助 +## 📚 术语解释 + +* **评测数据集**：用来测试 GPU 性能的一组标准化样本，包括代码、输入数据和预期结果。 + +* **GT（Ground Truth）**：标准参考答案或结果，用来验证程序运行是否正确。 + +* **吞吐量（Throughput）**：每秒钟能处理的数据量，越高表示 GPU 处理能力越强。 + +* **内存带宽（Memory Bandwidth）**：单位时间内 GPU 内存与计算核心之间的数据传输速度。 + +* **Prompt**：引导大语言模型（LLM）生成代码或内容的提示词。 + +* **LLM**：Large Language Model，大语言模型，如 ChatGPT、LLaMA 等。 + + +--- + +## 📬 联系与帮助 + +如需更多信息或格式说明，请查看官方文档或在本仓库提交[想法](https://gitee.com/ccf-ai-infra/GPUCodeForces/issues/new?template=feature.yml)进行讨论。  祝你挑战成功，贡献出高质量的 GPU 评测数据集！🚀 + +## FAQ -如需更多信息或格式说明，请查看官方文档或在本仓库提交 Issue 进行讨论。 -祝你挑战成功，贡献出高质量的 GPU 评测数据集！🚀 +[第一季FAQ参考](FAQ.md) \ No newline at end of file diff --git a/example/002-example/example_cudacode.py b/example/002-example/example_cudacode.py new file mode 100644 index 0000000000000000000000000000000000000000..3232300efee7e8feeb1170f2ea03d12b1295c238 --- /dev/null +++ b/example/002-example/example_cudacode.py @@ -0,0 +1,49 @@ +import torch +import torch.nn as nn +from torch.utils.cpp_extension import load_inline + +# 更简单的实现：只优化ReLU部分，矩阵乘法使用PyTorch +relu_source = """ +#include +#include + +__global__ void relu_kernel(const float* x, float* y, int size) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < size) { + y[idx] = fmaxf(x[idx], 0.f); + } +} + +torch::Tensor relu_cuda(torch::Tensor x) { + auto size = x.numel(); + auto y = torch::empty_like(x); + const int block_size = 256; + int num_blocks = (size + block_size - 1) / block_size; + relu_kernel<<>>(x.data_ptr(), y.data_ptr(), size); + return y; +} +""" + +relu_cpp_source = """ +torch::Tensor relu_cuda(torch::Tensor x); +""" + +# Compile the inline CUDA code +relu = load_inline( + name="relu", + cpp_sources=relu_cpp_source, + cuda_sources=relu_source, + functions=["relu_cuda"], + verbose=True +) + +class ModelNew(torch.nn.Module): + def __init__(self, weight): + super(ModelNew, self).__init__() + self.weight = nn.Parameter(weight) + self.relu = relu # The module containing the kernel + + def forward(self, x): + # 使用PyTorch的矩阵乘法，只优化ReLU部分 + x = torch.matmul(x, self.weight) + return self.relu.relu_cuda(x) diff --git a/example/002-example/example_torchcode.py b/example/002-example/example_torchcode.py new file mode 100644 index 0000000000000000000000000000000000000000..7e9d5f81ee131019d1d7b05c161a44ba01c2bf30 --- /dev/null +++ b/example/002-example/example_torchcode.py @@ -0,0 +1,35 @@ +import torch +import torch.nn as nn + +class Model(nn.Module): + """ + Model that performs matrix multiplication followed by ReLU activation. + """ + def __init__(self, weight): + super(Model, self).__init__() + self.weight = nn.Parameter(weight) + + def forward(self, x: torch.Tensor) -> torch.Tensor: + """ + Performs matrix multiplication and applies ReLU activation. + + Args: + x (torch.Tensor): Input tensor of shape [batch_size, input_dim] + + Returns: + torch.Tensor: Output tensor of shape [batch_size, output_dim] + """ + x = torch.matmul(x, self.weight) + return torch.relu(x) + +batch_size = 16 +input_dim = 1024 +output_dim = 2048 + +def get_inputs(): + x = torch.randn(batch_size, input_dim) + return [x] + +def get_init_inputs(): + weight = torch.randn(input_dim, output_dim) + return [weight] \ No newline at end of file diff --git a/example/002-example/prompt.txt b/example/002-example/prompt.txt new file mode 100644 index 0000000000000000000000000000000000000000..0deaedc34e51c38683026b099aefb08c820ea7f9 --- /dev/null +++ b/example/002-example/prompt.txt @@ -0,0 +1,30 @@ +You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups. + +You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination. + +Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is: + +```python +import torch +import torch.nn as nn +import torch.nn.functional as F + + +class Model(nn.Module): + def __init__(self) -> None: + super().__init__() + + def forward(self, a, b): + return a + b + + +def get_inputs(): + # randomly generate input tensors based on the model architecture + a = torch.randn(1, 128).cuda() + b = torch.randn(1, 128).cuda() + return [a, b] + + +def get_init_inputs(): + # randomly generate tensors required for initialization based on the model architecture + return [] \ No newline at end of file diff --git a/example/002-example/readme.md b/example/002-example/readme.md new file mode 100644 index 0000000000000000000000000000000000000000..1447c24c7d000100b76060b829fff2fcbb93515f --- /dev/null +++ b/example/002-example/readme.md @@ -0,0 +1,9 @@ +# 文件说明 + +example_torchcode.py:torch代码示例 + +example_cudacode.py：和torch对应的cuda代码 + +prompt.txt：利用LLM从torch代码生成cuda代码的prompt示例，（原始torch代码被附在prompt最后） + +run_code.py：用于测试生成的cuda代码和原始torch输出是否一致以及加速情况的示例代码 diff --git a/example/002-example/run_code.py b/example/002-example/run_code.py new file mode 100644 index 0000000000000000000000000000000000000000..24e869475896c03b5ea8ad44742706f3182d9154 --- /dev/null +++ b/example/002-example/run_code.py @@ -0,0 +1,88 @@ +########################################################### +# 性能和精度验证程序 +########################################################### +import torch +import torch.nn as nn +import time +from example_torchcode import Model, get_inputs, get_init_inputs +from example_cudacode import ModelNew + +def run_benchmark(): + # 检查 CUDA 是否可用 + if not torch.cuda.is_available(): + print("CUDA 不可用，请确保您有可用的 NVIDIA GPU 并已正确安装 PyTorch CUDA 版本。") + return + else: + device = torch.device("cuda") + + # 初始化模型 + init_inputs = get_init_inputs() + init_inputs = [ + x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in init_inputs + ] + inputs = get_inputs() + inputs = [ + x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in inputs + ] + + torch_model = Model(*init_inputs).cuda() + cuda_model = ModelNew(*init_inputs).cuda() + + torch_model.eval() + cuda_model.eval() + + print("-------------------- 精度对齐验证 --------------------") + with torch.no_grad(): + output_torch = torch_model(*inputs) + output_cuda = cuda_model(*inputs) + + # 更严格的精度检查 + abs_diff = (output_torch - output_cuda).abs() + max_diff = abs_diff.max().item() + mean_diff = abs_diff.mean().item() + + print(f"最大差异: {max_diff:.6f}") + print(f"平均差异: {mean_diff:.6f}") + + precision_flag = torch.allclose(output_torch, output_cuda, rtol=1e-05, atol=1e-05) + if precision_flag: + print("✅ 精度对齐：两个模型的输出结果非常接近。") + else: + print("❌ 精度不一致！") + + print("\n-------------------- 性能加速比测试 --------------------") + num_iterations = 1000 # 增加迭代次数以获得更准确的时间测量 + + # Warm up + for _ in range(100): + _ = torch_model(*inputs) + _ = cuda_model(*inputs) + + # PyTorch 模型计时 + torch.cuda.synchronize() + start_time = time.time() + for _ in range(num_iterations): + _ = torch_model(*inputs) + torch.cuda.synchronize() + torch_time = (time.time() - start_time) / num_iterations + + # 自定义 CUDA 内核计时 + torch.cuda.synchronize() + start_time = time.time() + for _ in range(num_iterations): + _ = cuda_model(*inputs) + torch.cuda.synchronize() + cuda_time = (time.time() - start_time) / num_iterations + + print(f"PyTorch (matmul + relu) 平均执行时间: {torch_time:.6f} 秒") + print(f"自定义 CUDA ReLU 平均执行时间: {cuda_time:.6f} 秒") + speedup = 0 + if cuda_time > 0: + speedup = torch_time / cuda_time + print(f"加速比 (Speedup): {speedup:.2f}x") + else: + print("CUDA 内核执行时间为0，无法计算加速比。") + return precision_flag, speedup + +if __name__ == "__main__": + precision_flag, speedup = run_benchmark() \ No newline at end of file diff --git "a/example/002-example/\345\217\202\350\265\233\350\200\205\351\234\200\350\246\201\346\217\220\344\276\233\347\232\204\345\206\205\345\256\271.md" "b/example/002-example/\345\217\202\350\265\233\350\200\205\351\234\200\350\246\201\346\217\220\344\276\233\347\232\204\345\206\205\345\256\271.md" new file mode 100644 index 0000000000000000000000000000000000000000..cb7588043339a54d759bed2514a78ca7014a9ef6 --- /dev/null +++ "b/example/002-example/\345\217\202\350\265\233\350\200\205\351\234\200\350\246\201\346\217\220\344\276\233\347\232\204\345\206\205\345\256\271.md" @@ -0,0 +1,4 @@ +1. 根据example_torchcode.py的格式，提供一个torch实现的op，命名为torchcode.py +2. 仿照prompt.txt的写法，利用llm（deepseek、通义千问、GPT、Gemini等大模型）生成一个初始的cuda算子，按照example_cudacode.py的格式组织成一个可以运行的cuda op，命名为cudacode_ori.py，并且利用run_code.py 检查算子精度 +3. 在符合精度要求的cudacode_ori.py基础上，进行cuda算子性能优化，用run_code.py检查算子精度和加速比，形成最终的最优性能的cuda算子实现，命名为cudacode_opt.py，格式符合example_cudacode.py +4. 针对每一个op，参赛者需要提供四个文件，torchcode.py、prompt.txt、cudacode_ori.py、example_cudacode.py \ No newline at end of file