Compare commits
No commits in common. "main" and "main" have entirely different histories.
327
README.md
327
README.md
|
|
@ -1,215 +1,181 @@
|
|||
# 春节不摆烂 · TileLang 算力入门挑战
|
||||
# 子赛题三、GPU 算子优化挑战赛
|
||||
|
||||
欢迎参加 **TileLang 算力入门挑战**！
|
||||
|
||||
本活动是沐曦股份联合 TileAI 技术社区推出的春节技术活动，旨在帮助 AI 开发者、学生及技术爱好者在春节期间入门 GPU 算子优化，掌握 TileLang 这一强大的深度学习编译器。
|
||||
欢迎参加 **GPU 算子优化挑战赛** 🎯！
|
||||
本比赛旨在通过优化深度学习框架中的核心计算模块，提升大模型的运行效率。在本仓库中，你可以提交你的优化代码、测试样例和使用说明。
|
||||
|
||||
---
|
||||
|
||||
## 活动简介
|
||||
## 🧠 比赛背景简介
|
||||
|
||||
春节，是技术宅一年中最容易摆烂、也最容易偷偷变强的时间。
|
||||
随着大语言模型（LLM, 如 ChatGPT 等）的广泛使用，其在运行推理时对计算资源的要求越来越高，出现了所谓的 “三高” 问题：
|
||||
|
||||
当别人刷短视频、抢红包的时候，总有人在凌晨 pull 模型、跑推理、调参数、看报错。
|
||||
- **高延迟**：响应慢
|
||||
- **高显存占用**：消耗大量显存
|
||||
- **高生态依赖**：对系统和软件依赖复杂
|
||||
|
||||
2026 年春节期间，沐曦股份联合 TileAI 技术社区，推出一次真正能上手、真有算力、真能学东西的春节技术活动：
|
||||
|
||||
> 不卷 KPI，不讲 PPT，就给你算力 + 课程 + 实操场景。
|
||||
本次挑战赛正是为了解决这些问题，鼓励选手**优化 GPU 上的底层算子（即最基础的数学运算模块）**，提高模型推理效率。
|
||||
|
||||
---
|
||||
|
||||
## 参与方式
|
||||
## 🔧 挑战方向
|
||||
|
||||
- 注册[沐曦开发者社区](https://external-sso.metax-tech.com/login?app_id=6840fb74b3b5d7665dcc0425&protocol=oidc&finish_login_url=%2Finteraction%2Foidc%2F475b38c5-7fb0-47e7-8258-16e945b83c79%2Flogin&login_page_context=)
|
||||
- 报名活动[领取算力券](https://developer.metax-tech.com/activities/2)
|
||||
本次挑战主要有两个技术方向：
|
||||
|
||||
### Test Time Scaling 算子优化
|
||||
- 对深度学习框架中的基础算子（如 PyTorch 或 PaddlePaddle 中的矩阵乘法）进行性能优化。
|
||||
- **GEMM（General Matrix Multiplication）** 是大模型中的核心计算操作。
|
||||
- 目标是对其内核（Kernel）进行精细调优，提高执行效率。
|
||||
|
||||
👉 **项目目标：在不更换硬件的前提下，让推理速度提升 30%！**
|
||||
|
||||
---
|
||||
|
||||
## 活动时间
|
||||
## 🚀 快速上手
|
||||
|
||||
**2026 年春节假期**
|
||||
本竞赛旨在评估参赛者在GPU并行计算领域的算法优化能力。为了快速让参赛者进入比赛状态，我们提供了三个核心算法的高性能版本参考，供参赛选手不断优化性能：
|
||||
- **ReduceSum**: 高精度归约求和
|
||||
- **SortPair**: 键值对稳定排序
|
||||
- **TopkPair**: 键值对TopK选择
|
||||
|
||||
- Day 0：活动预热 & 报名
|
||||
- Day 1-2：新手快速上手（领算力 + 跑第一个任务）
|
||||
- Day 3-9：进阶实操挑战（课程 + 任务解锁）
|
||||
- Day 10：春节技术打卡收官 & 排行榜公布
|
||||
[三个核心算法赛题模板](https://gitlink.org.cn/ccf-ai-infra/GPUKernelContest/tree/main/cp_template)
|
||||
|
||||
每天 1-2 小时即可参与，不打扰过年。
|
||||
### 📥 选手赛题准备
|
||||
|
||||
---
|
||||
1. 点击 **[创建赛题](https://gitlink.org.cn/ccf-ai-infra/GPUKernelContest/issues/new)** ，并记录赛题的ID
|
||||
2. 算力平台启动一个MACA3.0.0.4+PyTorch2.4.0的容器算力，详细步骤参考： **[算力平台使用说明](https://ai.gitee.com/docs/compute/container)**
|
||||
3. 用ssh或者vscode进入容器环境，Clone自己Fork的仓库
|
||||
```bash
|
||||
git clone https://gitlink.com/gitlinkuseid/GPUKernelContest.git
|
||||
```
|
||||
> 备注：`gitlinkuseid` 替换为您的URL。
|
||||
4. 进入参赛项目，创建以赛题ID为名称的目录。例如：[赛题3](https://gitlink.org.cn/ccf-ai-infra/GPUKernelContest/issues/3)
|
||||
```bash
|
||||
# 进入Clone的仓库
|
||||
cd GPUKernelContest
|
||||
# 创建以赛题ID为名称的目录
|
||||
mkdir -p S1/3
|
||||
```
|
||||
```
|
||||
# 创建后的目录结构如下:
|
||||
GPUKernelContest
|
||||
├── cp_template(说明：赛目模板目录)
|
||||
├── S1(说明：第一季比赛名)
|
||||
│ ├── 3(说明：以自己创建赛题ID命名目录存放自己需要提交的内容)
|
||||
```
|
||||
4. Fork仓库并初始化比赛环境（三个核心算法题优化赛题以外自定义的赛题需有入口run.sh脚本，供CI自动测试验证）
|
||||
1. 拷贝赛题样例`cp_template`到赛题`3`目录
|
||||
```bash
|
||||
# cp -r cp_template/* S1/3
|
||||
```
|
||||
2. 拷贝后的目录结构如下：
|
||||
```
|
||||
├── S1(说明：第一季比赛名)
|
||||
│ ├── 3(说明：以自己创建赛题ID命名目录存放自己需要提交的内容)
|
||||
| | ├── utils
|
||||
│ | ├── reduce_sum_algorithm.maca
|
||||
│ | ├── run.sh（说明：作为CI自动测试验证的入口）
|
||||
| | └── sort_pair_algorithm.maca
|
||||
| | └── topk_pair_algorithm.maca
|
||||
│ └── ……
|
||||
```
|
||||
|
||||
## 参与对象
|
||||
### 编译和测试
|
||||
|
||||
- AI / ML / Infra 工程师
|
||||
- 在校学生（计算机 / 人工智能相关专业）
|
||||
- 技术宅 / 开源爱好者
|
||||
- 想趁春节低成本入门 AI Infra / 大模型的新人
|
||||
|
||||
**不要求基础很强，但要求：愿意动手。**
|
||||
|
||||
---
|
||||
|
||||
## 三大任务方向
|
||||
|
||||
本活动主要分为三个任务方向，完成任意方向即可获得算力奖励：
|
||||
|
||||
### 任务一：mcTileLang 源码安装
|
||||
|
||||
完成 mcTileLang 源码编译安装，理解 TileLang 核心架构。
|
||||
|
||||
**目标：**
|
||||
- 成功编译安装 mcTileLang
|
||||
- 运行示例代码验证安装
|
||||
- 提交安装过程中遇到的问题与解决方案（如果有）
|
||||
|
||||
**详细教程：** [mcTileLang 源码编译教程](./spring_tilelang/docs/mctilelang-install.md)
|
||||
|
||||
---
|
||||
|
||||
### 任务二：TileLang Puzzle 练习
|
||||
|
||||
通过 TileLang Puzzle 练习，掌握 TileLang 编程基础。
|
||||
|
||||
**目标：**
|
||||
- 完成基础 Puzzle 任务
|
||||
- 理解 TileLang 的编程模型
|
||||
|
||||
**练习路径：** [tilelang-puzzles](https://github.com/tile-ai/tilelang-puzzles)
|
||||
|
||||
|
||||
---
|
||||
|
||||
### 任务三：Puzzle 学习文档撰写
|
||||
|
||||
整理学习笔记，编写 TileLang 学习文档，分享技术心得。
|
||||
|
||||
**目标：**
|
||||
- 整理 Puzzle 练习过程中的笔记
|
||||
- 编写 TileLang 学习文档
|
||||
- 分享技术心得与踩坑经历
|
||||
|
||||
**提交方式：** 将项目提交到 `spring_tilelang/submissions/YOUR_ID/` 目录
|
||||
|
||||
---
|
||||
|
||||
## 算力领取方式
|
||||
|
||||
### 基础算力包
|
||||
|
||||
报名成功即可领取春节专属基础算力包（100元福利）。
|
||||
|
||||
### 进阶算力包
|
||||
|
||||
完成指定任务可解锁：
|
||||
- 进阶算力包（春节专属：200元福利）
|
||||
- 高性能 GPU 限时体验
|
||||
|
||||
> 算力不是抽奖，是用行动换的。
|
||||
|
||||
---
|
||||
|
||||
## 任务等级
|
||||
|
||||
### 基础任务
|
||||
|
||||
| 等级 | 内容描述 |
|
||||
|------|----------|
|
||||
| Level 1 | 完成 mcTileLang 源码安装并提交验证截图 |
|
||||
| Level 2 | 完成 5 个及以上 TileLang Puzzle 练习 |
|
||||
| Level 3 | 提交 TileLang 学习文档 / 每个 Puzzle 一个文档或者学习总结 |
|
||||
|
||||
|
||||
## 奖励与福利
|
||||
|
||||
### 基础福利（限前50名）
|
||||
|
||||
- 春节专属算力包
|
||||
|
||||
### 进阶奖励（只要你敢，管够）
|
||||
|
||||
- 高性能 GPU 算力体验 + TileLang 进阶探索
|
||||
|
||||
### 彩蛋奖励（Top 5 活跃参与者）
|
||||
|
||||
- 技术周边（书籍、周边礼物等）
|
||||
- 优秀技术分享推荐至社区首页
|
||||
|
||||
---
|
||||
|
||||
## 快速开始
|
||||
|
||||
### 1. Fork 本仓库
|
||||
|
||||
点击右上角 Fork 按钮，将本仓库 Fork 到你的账号下。
|
||||
|
||||
### 2. Clone 你的仓库
|
||||
选手赛题目录内提供了编译、测试的脚本，供选手熟悉比赛环境，步骤如下：
|
||||
|
||||
```bash
|
||||
git clone https://gitlink.org.cn/YOUR_USERNAME/GPUKernelContest.git
|
||||
cd GPUKernelContest/spring_tilelang
|
||||
# ！！！注意参赛选手需要根据自己的赛题ID进入自己完成题目的目录！！！！
|
||||
cd S1/3
|
||||
```
|
||||
|
||||
### 3. 创建你的任务目录
|
||||
#### 1. 编译和运行
|
||||
|
||||
编译并运行所有算法测试（默认行为），如下：
|
||||
```bash
|
||||
./run.sh
|
||||
```
|
||||
|
||||
单个或几个赛题测试验证，修改`run.sh`脚本，详细如下：
|
||||
```bash
|
||||
#!/bin/bash
|
||||
|
||||
# 单个赛题测试验证(ReduceSum算法)
|
||||
./build_and_run.sh --run_reduce
|
||||
```
|
||||
|
||||
编译运行单个ReduceSum测试如下：
|
||||
```bash
|
||||
./run.sh # ReduceSum算法
|
||||
```
|
||||
|
||||
#### 2. 手动运行测试
|
||||
|
||||
```bash
|
||||
# 以你的 Gitlink ID 创建目录
|
||||
mkdir -p spring_tilelang/submissions/YOUR_ID
|
||||
cd spring_tilelang/submissions/YOUR_ID
|
||||
# 仅编译所有算法，不运行测试
|
||||
./build_and_run.sh --build-only
|
||||
|
||||
# 单个运行不同算法的测试
|
||||
./build/test_reducesum [correctness|performance|all]
|
||||
./build/test_sortpair [correctness|performance|all]
|
||||
./build/test_topkpair [correctness|performance|all]
|
||||
```
|
||||
|
||||
### 4. 开始任务
|
||||
对于如何提交可参考：[如何贡献](https://gitlink.org.cn/ccf-ai-infra/GPUKernelContest/tree/main/how-to-contribute.md)
|
||||
|
||||
根据上述三大任务方向，选择你感兴趣的任务开始：
|
||||
### ✅ 参赛要求：
|
||||
- 提交内容必须可以在MACA软件上运行。
|
||||
- 所提交的优化代码将由主办方审核，**需成功合并（Merge）到赛事官方仓库，才算有效提交。**
|
||||
|
||||
- [任务一：mcTileLang 源码安装](./spring_tilelang/docs/mctilelang-install.md)
|
||||
- [任务二：TileLang Puzzle 练习](./spring_tilelang/docs/tilelang-puzzle.md)
|
||||
- [任务三：学习文档撰写](./spring_tilelang/docs/writing-guide.md)
|
||||
### 📦 提交内容包含：
|
||||
- 算子优化后的代码
|
||||
- 可运行的测试用例
|
||||
- 使用说明文档
|
||||
|
||||
---
|
||||
|
||||
## 提交规范
|
||||
## 📈 评分机制
|
||||
|
||||
### 提交内容包含
|
||||
每次合并的提交会按以下规则评分，[mcTileLang](https://gitee.com/metax-maca/mcTileLang)详见[mcTileLang](docs/Tilelang/TileLang.md)：
|
||||
|
||||
- 任务完成代码 / 文档
|
||||
- 运行截图或验证结果
|
||||
- 使用说明文档（README.md）
|
||||
### 🎯 基础得分（Level）：
|
||||
| 等级 | 内容描述 | 分值 |
|
||||
|------|----------|------|
|
||||
| Level 1 | 优化一个 PyTorch或Paddle 算子 / 验证[mcTileLang](https://gitee.com/metax-maca/mcTileLang)的docs文件夹下的文档并提交验证结果的截图到对应的issue | 5 分 |
|
||||
| Level 2 | 融合优化 2~9 个算子 / 迁移[mcTileLang](https://gitee.com/metax-maca/mcTileLang)的docs文件夹下的文档已有的与CUDA相关的文档到MACA，并给对应的文档提交PR | 10 分 |
|
||||
| Level 3 | 为[mcTileLang](https://gitee.com/metax-maca/mcTileLang)的docs/deeplearning_operators文件夹下未编写文档的算子提交算子解读的PR/补充和修复已有文档 | 20 分 |
|
||||
| Level 4 | 含 MMA（多维矩阵乘）融合算子 / 用于大模型推理的复杂融合算子 / 给开源仓库[mcTileLang](https://gitee.com/metax-maca/mcTileLang)提交example文件夹下的代码PR | 50 分 |
|
||||
| 合并至MACA开源项目仓库的每个PR | 参考：[mcTVM](https://github.com/metax-maca/mcTVM),[mcTileLang](https://gitee.com/metax-maca/mcTileLang) | 50 分 |
|
||||
|
||||
### 提交方式
|
||||
> 注释事项，非AI Infra组下的项目PR需在赛题Issue中提供合并记录，并确保和参赛时使用邮箱一致的提交邮箱方为有效。
|
||||
|
||||
1. 在你的 Fork 仓库中完成修改
|
||||
2. 提交 Pull Request 到本仓库
|
||||
3. PR 标题格式：`[SpringTileLang] Gitlink ID - 任务名称`
|
||||
4. PR 描述中说明完成的任务内容和得分项
|
||||
### ✨ 加分项：
|
||||
| 内容 | 分值 |
|
||||
|------|------|
|
||||
| 代码规范、清晰 | +10 分 |
|
||||
| 性能优化明显 | +10 分 |
|
||||
| 记录优化过程、说明模型来源 | +20 分 |
|
||||
| 使用 LLM Prompt 自动生成代码及样例 | +20 分 |
|
||||
|
||||
### 提交示例
|
||||
|
||||
```
|
||||
[SpringTileLang] alice - mcTileLang 源码安装 + 5个Puzzle
|
||||
|
||||
完成任务：
|
||||
- [x] mcTileLang 源码安装（Level 1）
|
||||
- [x] TileLang Puzzle 练习（Level 2）
|
||||
- copy
|
||||
- conv
|
||||
```
|
||||
**注意：** 加分项只针对于在基础得分相同的情况下通过加分项来区分不同的排名和后续优秀选手的参考。
|
||||
|
||||
---
|
||||
|
||||
## 参考资源
|
||||
## 🏆 排名机制
|
||||
|
||||
### mcTileLang 相关
|
||||
1. 评委评分从高到低排序
|
||||
2. **评估规则：** 取前 12 名作为最终获奖选手
|
||||
3. 若基础得分相同：
|
||||
- 加分项多者优先
|
||||
- 提交数量多者优先
|
||||
- 提交时间早者优先
|
||||
4. 当同一参赛选手在本赛题有多个赛题的提交时，多个赛题计算累计得分
|
||||
|
||||
- [mcTileLang 仓库](https://gitee.com/metax-maca/mcTileLang)
|
||||
- [mcTileLang 文档](https://gitee.com/metax-maca/mcTileLang/tree/main/docs)
|
||||
- [TileLang 官方仓库](https://github.com/tile-ai/tilelang)
|
||||
---
|
||||
|
||||
### TileLang Puzzle
|
||||
|
||||
- [tilelang-puzzles 仓库](https://github.com/tile-ai/tilelang-puzzles)
|
||||
|
||||
### MACA 开源项目
|
||||
## 📚 参考MACA开源项目仓库
|
||||
|
||||
你可以参考以下项目仓库，了解算子开发与提交格式。如果为[GitHub](https://github.com/orgs/MetaX-MACA/repositories)、[Gitee](https://gitee.com/organizations/metax-maca/projects)仓库里面的项目提出一个好的Issue可以获得算力券的激励。
|
||||
如：
|
||||
- [mcTVM](https://github.com/MetaX-MACA/mcTVM)
|
||||
- [FlashMLA](https://github.com/MetaX-MACA/FlashMLA)
|
||||
- [mcEigen](https://github.com/MetaX-MACA/mcEigen)
|
||||
|
|
@ -217,25 +183,20 @@ cd spring_tilelang/submissions/YOUR_ID
|
|||
|
||||
---
|
||||
|
||||
## 💡 术语解释
|
||||
|
||||
|
||||
## 联系与帮助
|
||||
|
||||
如需帮助或有疑问，请：
|
||||
|
||||
1. 在本仓库发起 Issue
|
||||
2. 加入 TileLang 微信群
|
||||
|
||||
<img src="./tilelang_group.png" alt="TileLang群" width="300" height="300">
|
||||
|
||||
3. 联系活动主办方
|
||||
- **算子（Operator）**：指深度学习框架中的基本计算模块，例如矩阵乘法、卷积等。
|
||||
- **GEMM**：全称 General Matrix Multiplication，一种用于矩阵计算的核心算法，是大模型中的基础运算。
|
||||
- **MMA**：Matrix Multiply-Accumulate，多维矩阵乘加运算，适用于复杂计算加速。
|
||||
- **LLM**：Large Language Model，大语言模型，如 GPT、BERT 等。
|
||||
- **推理（Inference）**：模型训练完成后，用来“预测”或“使用”的过程。
|
||||
- **Prompt**：用于引导大模型生成特定内容的输入提示词。
|
||||
- **PR（Pull Request）**：在 Git 仓库中提交你的修改请求，供维护者审查后合并。
|
||||
|
||||
---
|
||||
|
||||
## 活动总结
|
||||
## 📬 联系与帮助
|
||||
|
||||
> **这是一次不讲虚话的春节技术活动，**
|
||||
> **算力是真的，课程是能跑的，**
|
||||
> **参与的人，是真的能学到东西的。**
|
||||
如需帮助或有疑问，请联系主办方或在项目中发起 Issue。
|
||||
|
||||
祝你春节不摆烂，用算力换成长！
|
||||
祝你挑战顺利，优化出更快的大模型推理体验！🚀
|
||||
|
|
|
|||
|
|
@ -0,0 +1,274 @@
|
|||
#!/bin/bash
|
||||
|
||||
# GPU高性能并行计算算法优化竞赛 - 统一编译和运行脚本
|
||||
# 整合了所有算法的编译、运行和公共配置
|
||||
|
||||
# ============================================================================
|
||||
# 公共配置和工具函数
|
||||
# ============================================================================
|
||||
|
||||
# 设置颜色
|
||||
RED='\033[0;31m'
|
||||
GREEN='\033[0;32m'
|
||||
BLUE='\033[0;34m'
|
||||
YELLOW='\033[0;33m'
|
||||
NC='\033[0m' # No Color
|
||||
|
||||
# 打印函数
|
||||
print_info() {
|
||||
echo -e "${BLUE}[INFO]${NC} $1"
|
||||
}
|
||||
|
||||
print_success() {
|
||||
echo -e "${GREEN}[SUCCESS]${NC} $1"
|
||||
}
|
||||
|
||||
print_error() {
|
||||
echo -e "${RED}[ERROR]${NC} $1"
|
||||
}
|
||||
|
||||
print_warning() {
|
||||
echo -e "${YELLOW}[WARNING]${NC} $1"
|
||||
}
|
||||
|
||||
# 编译配置 - 可通过环境变量自定义
|
||||
COMPILER=${COMPILER:-mxcc}
|
||||
COMPILER_FLAGS=${COMPILER_FLAGS:-"-O3 -std=c++17 --extended-lambda -DRUN_FULL_TEST"}
|
||||
|
||||
# ***** 这里是关键修改点1：头文件目录 *****
|
||||
# 现在头文件在 utils/ 目录下
|
||||
HEADER_DIR=${HEADER_DIR:-utils}
|
||||
|
||||
# ***** 这里是关键修改点2：源文件目录 *****
|
||||
# 现在源文件在 ./ 目录下
|
||||
SOURCE_CODE_DIR=${SOURCE_CODE_DIR:-}
|
||||
|
||||
BUILD_DIR=${BUILD_DIR:-build}
|
||||
|
||||
# 编译单个算法的通用函数
|
||||
# 参数: $1=算法名称, $2=源文件名（不含路径）
|
||||
compile_algorithm() {
|
||||
local algo_name="$1"
|
||||
local source_file_name="$2" # 例如 "reduce_sum_algorithm.maca"
|
||||
local target_file="$BUILD_DIR/test_${algo_name,,}" # 转换为小写
|
||||
|
||||
print_info "编译 $algo_name 算法..."
|
||||
|
||||
# 创建构建目录
|
||||
mkdir -p "$BUILD_DIR"
|
||||
|
||||
# ***** 这里是关键修改点3：编译命令 *****
|
||||
# -I$HEADER_DIR 用于告诉编译器头文件在哪里
|
||||
# $SOURCE_CODE_DIR/$source_file_name 用于指定要编译的源文件的完整路径
|
||||
local compile_cmd="$COMPILER $COMPILER_FLAGS -I$HEADER_DIR $source_file_name -o $target_file"
|
||||
|
||||
print_info "执行: $compile_cmd"
|
||||
|
||||
if $compile_cmd; then
|
||||
print_success "$algo_name 编译完成!"
|
||||
echo ""
|
||||
echo "运行测试:"
|
||||
echo " ./$target_file [correctness|performance|all]"
|
||||
return 0
|
||||
else
|
||||
print_error "$algo_name 编译失败!"
|
||||
return 1
|
||||
fi
|
||||
}
|
||||
|
||||
# 显示编译配置信息
|
||||
show_build_config() {
|
||||
print_info "编译配置:"
|
||||
echo " COMPILER: $COMPILER"
|
||||
echo " COMPILER_FLAGS: $COMPILER_FLAGS"
|
||||
echo " HEADER_DIR: $HEADER_DIR" # 显示头文件目录
|
||||
echo " SOURCE_CODE_DIR: $SOURCE_CODE_DIR" # 显示源文件目录
|
||||
echo " BUILD_DIR: $BUILD_DIR"
|
||||
echo ""
|
||||
}
|
||||
|
||||
# 运行单个测试
|
||||
run_single_test() {
|
||||
local algo_name="$1"
|
||||
local test_mode="${2:-all}"
|
||||
local test_file="$BUILD_DIR/test_${algo_name,,}"
|
||||
|
||||
if [ -f "$test_file" ]; then
|
||||
print_info "运行 $algo_name 测试 (模式: $test_mode)..."
|
||||
"./$test_file" "$test_mode"
|
||||
return $?
|
||||
else
|
||||
print_error "$algo_name 测试程序不存在: $test_file"
|
||||
return 1
|
||||
fi
|
||||
}
|
||||
|
||||
# ============================================================================
|
||||
# 主脚本逻辑
|
||||
# ============================================================================
|
||||
|
||||
# 显示帮助信息 (整合了所有选项)
|
||||
show_help() {
|
||||
echo "GPU算法竞赛统一编译和运行脚本"
|
||||
echo "用法: $0 [选项]"
|
||||
echo ""
|
||||
echo "选项:"
|
||||
echo " --help 显示帮助信息"
|
||||
echo " --build-only 仅编译所有算法，不运行测试"
|
||||
echo " --run_reduce [MODE] 编译并运行ReduceSum算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo " --run_sort [MODE] 编译并运行SortPair算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo " --run_topk [MODE] 编译并运行TopkPair算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo ""
|
||||
echo "示例:"
|
||||
echo " $0 # 编译并运行所有测试（默认行为）"
|
||||
echo " $0 --build-only # 仅编译所有算法"
|
||||
echo " $0 --run_sort performance # 编译并运行SortPair性能测试"
|
||||
echo ""
|
||||
}
|
||||
|
||||
# 解析命令行参数
|
||||
RUN_MODE="run_all" # 默认为编译并运行所有测试
|
||||
ALGO_TO_RUN="" # 记录要运行的单个算法
|
||||
SINGLE_ALGO_TEST_MODE="all" # 单个算法的测试模式
|
||||
|
||||
while [[ $# -gt 0 ]]; do
|
||||
case $1 in
|
||||
--help)
|
||||
show_help
|
||||
exit 0
|
||||
;;
|
||||
--build-only)
|
||||
RUN_MODE="build_only"
|
||||
shift
|
||||
;;
|
||||
--run_reduce)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="ReduceSum"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
--run_sort)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="SortPair"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
--run_topk)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="TopkPair"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
*)
|
||||
print_error "未知选项: $1"
|
||||
show_help
|
||||
exit 1
|
||||
;;
|
||||
esac
|
||||
done
|
||||
|
||||
if [ "$RUN_MODE" = "build_only" ]; then
|
||||
print_info "开始编译所有算法..."
|
||||
else
|
||||
print_info "开始编译并运行所有算法..."
|
||||
fi
|
||||
print_info "工作目录: $(pwd)"
|
||||
print_info "编译时间: $(date '+%Y-%m-%d %H:%M:%S')"
|
||||
show_build_config
|
||||
|
||||
# 清理构建目录
|
||||
if [ -d "$BUILD_DIR" ]; then
|
||||
print_info "清理现有构建目录: $BUILD_DIR"
|
||||
rm -rf "$BUILD_DIR"
|
||||
fi
|
||||
|
||||
# 核心逻辑：根据 RUN_MODE 执行操作
|
||||
case "$RUN_MODE" in
|
||||
"build_only")
|
||||
print_info "编译所有算法..."
|
||||
|
||||
# 直接调用 compile_algorithm 函数
|
||||
print_info "[1/3] 编译ReduceSum..."
|
||||
if ! compile_algorithm "ReduceSum" "reduce_sum_algorithm.maca"; then
|
||||
print_error "ReduceSum编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[2/3] 编译SortPair..."
|
||||
if ! compile_algorithm "SortPair" "sort_pair_algorithm.maca"; then
|
||||
print_error "SortPair编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[3/3] 编译TopkPair..."
|
||||
if ! compile_algorithm "TopkPair" "topk_pair_algorithm.maca"; then
|
||||
print_error "TopkPair编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_success "所有算法编译完成!"
|
||||
echo ""
|
||||
echo "可执行文件:"
|
||||
echo " $BUILD_DIR/test_reducesum - ReduceSum算法测试"
|
||||
echo " $BUILD_DIR/test_sortpair - SortPair算法测试"
|
||||
echo " $BUILD_DIR/test_topkpair - TopkPair算法测试"
|
||||
echo ""
|
||||
echo "使用方法:"
|
||||
echo " ./$BUILD_DIR/test_reducesum [correctness|performance|all]"
|
||||
echo " ./$BUILD_DIR/test_sortpair [correctness|performance|all]"
|
||||
echo " ./$BUILD_DIR/test_topkpair [correctness|performance|all]"
|
||||
;;
|
||||
|
||||
"run_all")
|
||||
print_info "编译并运行所有算法测试..."
|
||||
|
||||
# 直接调用 compile_algorithm 和 run_single_test 函数
|
||||
print_info "[1/3] ReduceSum..."
|
||||
if compile_algorithm "ReduceSum" "reduce_sum_algorithm.maca"; then
|
||||
run_single_test "ReduceSum" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[2/3] SortPair..."
|
||||
if compile_algorithm "SortPair" "sort_pair_algorithm.maca"; then
|
||||
run_single_test "SortPair" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[3/3] TopkPair..."
|
||||
if compile_algorithm "TopkPair" "topk_pair_algorithm.maca"; then
|
||||
run_single_test "TopkPair" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_success "所有测试完成!"
|
||||
;;
|
||||
|
||||
"run_single")
|
||||
print_info "编译并运行 ${ALGO_TO_RUN} 测试 (模式: ${SINGLE_ALGO_TEST_MODE})..."
|
||||
local source_file_name=""
|
||||
case "$ALGO_TO_RUN" in
|
||||
"ReduceSum") source_file_name="reduce_sum_algorithm.maca" ;;
|
||||
"SortPair") source_file_name="sort_pair_algorithm.maca" ;;
|
||||
"TopkPair") source_file_name="topk_pair_algorithm.maca" ;;
|
||||
esac
|
||||
|
||||
if compile_algorithm "$ALGO_TO_RUN" "$source_file_name"; then
|
||||
run_single_test "$ALGO_TO_RUN" "$SINGLE_ALGO_TEST_MODE"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
;;
|
||||
esac
|
||||
|
|
@ -0,0 +1,97 @@
|
|||
# 样例赛题说明
|
||||
|
||||
## GPU高性能并行计算算法优化
|
||||
|
||||
要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。
|
||||
|
||||
在正确性、稳定性前提下，比拼算法性能。
|
||||
|
||||
# 1. ReduceSum算法优化
|
||||
```cpp
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
class ReduceSumAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：InputT: float, OutputT: float
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入d\_in数据量为num\_items
|
||||
|
||||
注意事项
|
||||
|
||||
* 累计误差不大于cpu double golden基准的0.5%
|
||||
* 注意针对NAN和INF等异常值的处理
|
||||
|
||||
|
||||
加分项
|
||||
|
||||
* 使用tensor core计算reduce
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
||||
# 2. Sort Pair算法优化
|
||||
```cpp
|
||||
template <typename KeyType, typename ValueType>
|
||||
class SortPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void sort(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, bool descending) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：key: float, value: int32\_t
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入、输出的key和value的数据量一致，均为num\_items
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 需要校验结果正确性
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的排序，如half、double、int32_t等
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
# 3. Topk Pair算法优化
|
||||
```cpp
|
||||
template <typename KeyType, typename ValueType>
|
||||
class TopkPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void topk(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, int k, bool descending) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：key: float, value: int32\_t
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入的key和value的数据量一致，为num\_items；输出的key和value的数据量一致，为k
|
||||
* k的范围：32，50，100，256，1024。k不大于num\_items
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的键值对，实现类型通用算法
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
|
@ -0,0 +1,277 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/reduce.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/functional.h>
|
||||
#endif
|
||||
|
||||
// 误差容忍度
|
||||
constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5%
|
||||
|
||||
// ============================================================================
|
||||
// ReduceSum算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
class ReduceSumAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能归约算法
|
||||
|
||||
// 示例：参赛者可以调用1个或多个自定义kernel
|
||||
// blockReduceKernel<<<grid, block>>>(d_in, temp_results, num_items, init_value);
|
||||
// finalReduceKernel<<<1, block>>>(temp_results, d_out, grid.x);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
auto input_ptr = thrust::device_pointer_cast(d_in);
|
||||
auto output_ptr = thrust::device_pointer_cast(d_out);
|
||||
|
||||
// 直接使用thrust::reduce进行归约
|
||||
*output_ptr = thrust::reduce(
|
||||
thrust::device,
|
||||
input_ptr,
|
||||
input_ptr + num_items,
|
||||
static_cast<OutputT>(init_value)
|
||||
);
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：中间结果缓冲区、多阶段归约等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "ReduceSum 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
ReduceSumAlgorithm<float, float> algorithm;
|
||||
|
||||
bool allPassed = true;
|
||||
|
||||
// 测试不同数据规模
|
||||
for (int i = 0; i < NUM_TEST_SIZES && i < 2; i++) { // 限制测试规模
|
||||
int size = std::min(TEST_SIZES[i], 10000);
|
||||
std::cout << " 测试规模: " << size << std::endl;
|
||||
|
||||
// 测试普通数据
|
||||
{
|
||||
auto data = generator.generateRandomFloats(size, -10.0f, 10.0f);
|
||||
float init_value = 1.0f;
|
||||
|
||||
// CPU参考计算
|
||||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||||
|
||||
// GPU计算
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证误差
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||||
std::cout << " 失败: 误差过大 " << relative_error << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过 (误差: " << relative_error << ")" << std::endl;
|
||||
}
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
|
||||
// 测试特殊值 (NaN, Inf)
|
||||
if (size > 100) {
|
||||
std::cout << " 测试特殊值..." << std::endl;
|
||||
auto data = generator.generateSpecialFloats(size);
|
||||
float init_value = 0.0f;
|
||||
|
||||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||||
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
// 对于包含特殊值的情况，检查是否正确处理
|
||||
if (std::isfinite(cpu_result) && std::isfinite(gpu_result)) {
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||||
std::cout << " 失败: 特殊值处理错误" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过 (特殊值处理)" << std::endl;
|
||||
}
|
||||
} else {
|
||||
std::cout << " 通过 (特殊值结果)" << std::endl;
|
||||
}
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
}
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
PerformanceDisplay::printReduceSumHeader();
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
ReduceSumAlgorithm<float, float> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
// 生成测试数据
|
||||
auto data = generator.generateRandomFloats(size);
|
||||
float init_value = 0.0f;
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
|
||||
// 计算性能指标
|
||||
auto metrics = PerformanceCalculator::calculateReduceSum(size, avg_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printReduceSumData(size, avg_time, metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["time_ms"] = std::to_string(avg_time);
|
||||
entry["throughput_gps"] = std::to_string(metrics.throughput_gps);
|
||||
entry["data_type"] = "float";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateReduceSumYAML(perf_data, "reduce_sum_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("reduce_sum_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== ReduceSum 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << ReduceSumAlgorithm<float, float>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,13 @@
|
|||
#!/bin/bash
|
||||
|
||||
# 单个赛题测试验证(ReduceSum算法)
|
||||
#./build_and_run.sh --run_reduce
|
||||
|
||||
# 单个赛题测试验证(SortPair算法)
|
||||
#./build_and_run.sh --run_reduce
|
||||
|
||||
# 单个赛题测试验证(TopkPair算法)
|
||||
# ./build_and_run.sh --run_topk
|
||||
|
||||
# 默认全量赛题测试验证，参赛选手单个优化，参考单个脚本执行方式，CI入口run.sh
|
||||
./build_and_run.sh
|
||||
|
|
@ -0,0 +1,275 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/iterator/zip_iterator.h>
|
||||
#include <thrust/tuple.h>
|
||||
#endif
|
||||
|
||||
// ============================================================================
|
||||
// SortPair算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename KeyType, typename ValueType>
|
||||
class SortPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void sort(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, bool descending) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能排序算法
|
||||
|
||||
// 示例：参赛者可以调用1个或多个自定义kernel
|
||||
// preprocessKernel<<<grid, block>>>(d_keys_in, d_values_in, num_items);
|
||||
// mainSortKernel<<<grid, block>>>(d_keys_out, d_values_out, num_items, descending);
|
||||
// postprocessKernel<<<grid, block>>>(d_keys_out, d_values_out, num_items);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_out, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_out, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
auto key_ptr = thrust::device_pointer_cast(d_keys_out);
|
||||
auto value_ptr = thrust::device_pointer_cast(d_values_out);
|
||||
|
||||
if (descending) {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：临时缓冲区、多个kernel函数、流等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "SortPair 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
SortPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
// 测试小规模数据
|
||||
int size = 10000;
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
// 测试升序和降序
|
||||
bool allPassed = true;
|
||||
for (bool descending : {false, true}) {
|
||||
std::cout << " " << (descending ? "降序" : "升序") << " 测试..." << std::endl;
|
||||
|
||||
// CPU参考结果
|
||||
auto cpu_keys = keys;
|
||||
auto cpu_values = values;
|
||||
cpuSortPair(cpu_keys, cpu_values, descending);
|
||||
|
||||
// GPU算法结果
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
|
||||
// 获取结果
|
||||
std::vector<float> gpu_keys(size);
|
||||
std::vector<uint32_t> gpu_values(size);
|
||||
MACA_CHECK(mcMemcpy(gpu_keys.data(), d_keys_out, size * sizeof(float), mcMemcpyDeviceToHost));
|
||||
MACA_CHECK(mcMemcpy(gpu_values.data(), d_values_out, size * sizeof(uint32_t), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证结果
|
||||
bool keysMatch = compareArrays(cpu_keys, gpu_keys, 1e-5);
|
||||
bool valuesMatch = compareArrays(cpu_values, gpu_values);
|
||||
|
||||
if (!keysMatch || !valuesMatch) {
|
||||
std::cout << " 失败: 结果不匹配" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_in);
|
||||
mcFree(d_values_out);
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
PerformanceDisplay::printSortPairHeader();
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
SortPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
// 生成测试数据
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
float asc_time = 0, desc_time = 0;
|
||||
|
||||
// 测试升序和降序
|
||||
for (bool descending : {false, true}) {
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
if (descending) {
|
||||
desc_time = avg_time;
|
||||
} else {
|
||||
asc_time = avg_time;
|
||||
}
|
||||
}
|
||||
|
||||
// 计算性能指标
|
||||
auto asc_metrics = PerformanceCalculator::calculateSortPair(size, asc_time);
|
||||
auto desc_metrics = PerformanceCalculator::calculateSortPair(size, desc_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printSortPairData(size, asc_time, desc_time, asc_metrics, desc_metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["asc_time_ms"] = std::to_string(asc_time);
|
||||
entry["desc_time_ms"] = std::to_string(desc_time);
|
||||
entry["asc_throughput_gps"] = std::to_string(asc_metrics.throughput_gps);
|
||||
entry["desc_throughput_gps"] = std::to_string(desc_metrics.throughput_gps);
|
||||
entry["key_type"] = "float";
|
||||
entry["value_type"] = "uint32_t";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_in);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateSortPairYAML(perf_data, "sort_pair_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("sort_pair_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== SortPair 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << SortPairAlgorithm<float, uint32_t>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,317 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
#include <fstream>
|
||||
#include <map>
|
||||
#include <chrono>
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/iterator/zip_iterator.h>
|
||||
#include <thrust/tuple.h>
|
||||
#include <thrust/copy.h>
|
||||
#endif
|
||||
|
||||
static const int TOPK_VALUES[] = {32, 50, 100, 256, 1024};
|
||||
static const int NUM_TOPK_VALUES = sizeof(TOPK_VALUES) / sizeof(TOPK_VALUES[0]);
|
||||
|
||||
// ============================================================================
|
||||
// TopkPair算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename KeyType, typename ValueType>
|
||||
class TopkPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void topk(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, int k, bool descending) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能TopK算法
|
||||
|
||||
// 示例：参赛者可以调用多个自定义kernel
|
||||
// TopkKernel1<<<grid, block>>>(d_keys_in, d_values_in, temp_results, num_items, k);
|
||||
// TopkKernel2<<<grid, block>>>(temp_results, d_keys_out, d_values_out, k, descending);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
|
||||
KeyType* temp_keys;
|
||||
ValueType* temp_values;
|
||||
MACA_CHECK(mcMalloc(&temp_keys, num_items * sizeof(KeyType)));
|
||||
MACA_CHECK(mcMalloc(&temp_values, num_items * sizeof(ValueType)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(temp_values, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
auto key_ptr = thrust::device_pointer_cast(temp_keys);
|
||||
auto value_ptr = thrust::device_pointer_cast(temp_values);
|
||||
|
||||
// 由于greater和less是不同类型，需要分别调用
|
||||
if (descending) {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_out, temp_keys, k * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_out, temp_values, k * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
mcFree(temp_keys);
|
||||
mcFree(temp_values);
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：分块大小、临时缓冲区、多流处理等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "TopkPair 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
TopkPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
int size = 10000;
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
bool allPassed = true;
|
||||
|
||||
// 测试不同k值
|
||||
for (int ki = 0; ki < NUM_TOPK_VALUES && ki < 4; ki++) { // 限制测试范围
|
||||
int k = TOPK_VALUES[ki];
|
||||
if (k > size) continue;
|
||||
|
||||
std::cout << " 测试 k=" << k << std::endl;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t)));
|
||||
|
||||
for (bool descending : {false, true}) {
|
||||
std::cout << " " << (descending ? "降序" : "升序") << " TopK..." << std::endl;
|
||||
|
||||
// CPU参考结果
|
||||
std::vector<float> cpu_keys_out;
|
||||
std::vector<uint32_t> cpu_values_out;
|
||||
cpuTopkPair(keys, values, cpu_keys_out, cpu_values_out, k, descending);
|
||||
|
||||
// GPU算法结果
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
|
||||
// 获取结果
|
||||
std::vector<float> gpu_keys_out(k);
|
||||
std::vector<uint32_t> gpu_values_out(k);
|
||||
MACA_CHECK(mcMemcpy(gpu_keys_out.data(), d_keys_out, k * sizeof(float), mcMemcpyDeviceToHost));
|
||||
MACA_CHECK(mcMemcpy(gpu_values_out.data(), d_values_out, k * sizeof(uint32_t), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证结果
|
||||
bool keysMatch = compareArrays(cpu_keys_out, gpu_keys_out, 1e-5);
|
||||
bool valuesMatch = compareArrays(cpu_values_out, gpu_values_out);
|
||||
|
||||
if (!keysMatch || !valuesMatch) {
|
||||
std::cout << " 失败: 结果不匹配" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_values_in);
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
std::cout << "\nTopkPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
TopkPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
// 针对不同数据规模测试
|
||||
for (int size_idx = 0; size_idx < NUM_TEST_SIZES; size_idx++) {
|
||||
int size = TEST_SIZES[size_idx];
|
||||
std::cout << "\n数据规模: " << size << std::endl;
|
||||
std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(74, '-') << std::endl;
|
||||
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in;
|
||||
uint32_t *d_values_in;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
for (int ki = 0; ki < NUM_TOPK_VALUES; ki++) {
|
||||
int k = TOPK_VALUES[ki];
|
||||
if (k > size) continue;
|
||||
|
||||
float *d_keys_out;
|
||||
uint32_t *d_values_out;
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t)));
|
||||
|
||||
float asc_time = 0, desc_time = 0;
|
||||
|
||||
for (bool descending : {false, true}) {
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
if (descending) {
|
||||
desc_time = avg_time;
|
||||
} else {
|
||||
asc_time = avg_time;
|
||||
}
|
||||
}
|
||||
|
||||
// 计算性能指标
|
||||
auto asc_metrics = PerformanceCalculator::calculateTopkPair(size, k, asc_time);
|
||||
auto desc_metrics = PerformanceCalculator::calculateTopkPair(size, k, desc_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printTopkPairData(k, asc_time, desc_time, asc_metrics, desc_metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["k_value"] = std::to_string(k);
|
||||
entry["asc_time_ms"] = std::to_string(asc_time);
|
||||
entry["desc_time_ms"] = std::to_string(desc_time);
|
||||
entry["asc_throughput_gps"] = std::to_string(asc_metrics.throughput_gps);
|
||||
entry["desc_throughput_gps"] = std::to_string(desc_metrics.throughput_gps);
|
||||
entry["key_type"] = "float";
|
||||
entry["value_type"] = "uint32_t";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_values_in);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateTopkPairYAML(perf_data, "topk_pair_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("topk_pair_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== TopkPair 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << TopkPairAlgorithm<float, uint32_t>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,114 @@
|
|||
#pragma once
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <string>
|
||||
|
||||
// ============================================================================
|
||||
// 性能计算和显示工具
|
||||
// ============================================================================
|
||||
|
||||
class PerformanceCalculator {
|
||||
public:
|
||||
// ReduceSum性能计算
|
||||
struct ReduceSumMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static ReduceSumMetrics calculateReduceSum(int size, float time_ms) {
|
||||
ReduceSumMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
|
||||
// SortPair性能计算
|
||||
struct SortPairMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static SortPairMetrics calculateSortPair(int size, float time_ms) {
|
||||
SortPairMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
|
||||
// TopkPair性能计算
|
||||
struct TopkPairMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static TopkPairMetrics calculateTopkPair(int size, int k, float time_ms) {
|
||||
TopkPairMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 性能显示工具
|
||||
// ============================================================================
|
||||
|
||||
class PerformanceDisplay {
|
||||
public:
|
||||
// 显示ReduceSum性能表头
|
||||
static void printReduceSumHeader() {
|
||||
std::cout << "\nReduceSum 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: float -> float" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
std::cout << std::setw(12) << "数据规模" << std::setw(15) << "时间(ms)"
|
||||
<< std::setw(20) << "吞吐量(G/s)" << std::endl;
|
||||
std::cout << std::string(47, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示SortPair性能表头
|
||||
static void printSortPairHeader() {
|
||||
std::cout << "\nSortPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
std::cout << std::setw(12) << "数据规模" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(78, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示TopkPair性能表头
|
||||
static void printTopkPairHeader() {
|
||||
std::cout << "\nTopkPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
}
|
||||
|
||||
static void printTopkPairDataHeader() {
|
||||
std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(74, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示性能数据行
|
||||
static void printReduceSumData(int size, float time_ms, const PerformanceCalculator::ReduceSumMetrics& metrics) {
|
||||
std::cout << std::setw(12) << size << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< time_ms << std::setw(20) << std::setprecision(3) << metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
static void printSortPairData(int size, float asc_time, float desc_time,
|
||||
const PerformanceCalculator::SortPairMetrics& asc_metrics,
|
||||
const PerformanceCalculator::SortPairMetrics& desc_metrics) {
|
||||
std::cout << std::setw(12) << size << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< asc_time << std::setw(15) << desc_time << std::setw(16) << std::setprecision(3)
|
||||
<< asc_metrics.throughput_gps << std::setw(16) << desc_metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
static void printTopkPairData(int k, float asc_time, float desc_time,
|
||||
const PerformanceCalculator::TopkPairMetrics& asc_metrics,
|
||||
const PerformanceCalculator::TopkPairMetrics& desc_metrics) {
|
||||
std::cout << std::setw(8) << k << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< asc_time << std::setw(15) << desc_time << std::setw(16) << std::setprecision(3)
|
||||
<< asc_metrics.throughput_gps << std::setw(16) << desc_metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
// 显示性能文件保存消息
|
||||
static void printSavedMessage(const std::string& filename) {
|
||||
std::cout << "\n性能结果已保存到: " << filename << std::endl;
|
||||
}
|
||||
};
|
||||
|
|
@ -0,0 +1,234 @@
|
|||
#pragma once
|
||||
#include <vector>
|
||||
#include <random>
|
||||
#include <algorithm>
|
||||
#include <mc_runtime.h>
|
||||
#include <maca_fp16.h>
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
#include <cmath>
|
||||
|
||||
// 引入模块化头文件
|
||||
#include "yaml_reporter.h"
|
||||
#include "performance_utils.h"
|
||||
|
||||
// ============================================================================
|
||||
// 测试配置常量
|
||||
// ============================================================================
|
||||
#ifndef RUN_FULL_TEST
|
||||
const int TEST_SIZES[] = {1000000, 134217728}; // 1M, 128M, 512M, 1G
|
||||
#else
|
||||
const int TEST_SIZES[] = {1000000, 134217728, 536870912, 1073741824}; // 1M, 128M, 512M, 1G
|
||||
#endif
|
||||
|
||||
const int NUM_TEST_SIZES = sizeof(TEST_SIZES) / sizeof(TEST_SIZES[0]);
|
||||
|
||||
// 性能测试重复次数
|
||||
constexpr int WARMUP_ITERATIONS = 5;
|
||||
constexpr int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
|
||||
// ============================================================================
|
||||
// 错误检查宏
|
||||
// ============================================================================
|
||||
#define MACA_CHECK(call) \
|
||||
do { \
|
||||
mcError_t error = call; \
|
||||
if (error != mcSuccess) { \
|
||||
std::cerr << "MACA error at " << __FILE__ << ":" << __LINE__ \
|
||||
<< " - " << mcGetErrorString(error) << std::endl; \
|
||||
exit(1); \
|
||||
} \
|
||||
} while(0)
|
||||
|
||||
// ============================================================================
|
||||
// 测试数据生成器
|
||||
// ============================================================================
|
||||
class TestDataGenerator {
|
||||
private:
|
||||
std::mt19937 rng;
|
||||
|
||||
public:
|
||||
TestDataGenerator(uint32_t seed = 42) : rng(seed) {}
|
||||
|
||||
// 生成随机float数组
|
||||
std::vector<float> generateRandomFloats(int size, float min_val = -1000.0f, float max_val = 1000.0f) {
|
||||
std::vector<float> data(size);
|
||||
std::uniform_real_distribution<float> dist(min_val, max_val);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = dist(rng);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机half数组
|
||||
std::vector<half> generateRandomHalfs(int size, float min_val = -100.0f, float max_val = 100.0f) {
|
||||
std::vector<half> data(size);
|
||||
std::uniform_real_distribution<float> dist(min_val, max_val);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = __float2half(dist(rng));
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机uint32_t数组
|
||||
std::vector<uint32_t> generateRandomUint32(int size) {
|
||||
std::vector<uint32_t> data(size);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = static_cast<uint32_t>(i); // 使用索引作为值，便于验证稳定排序
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机int64_t数组
|
||||
std::vector<int64_t> generateRandomInt64(int size) {
|
||||
std::vector<int64_t> data(size);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = static_cast<int64_t>(i);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成包含NaN和Inf的测试数据 (half版本)
|
||||
std::vector<half> generateSpecialHalfs(int size) {
|
||||
std::vector<half> data = generateRandomHalfs(size, -10.0f, 10.0f);
|
||||
if (size > 100) {
|
||||
data[10] = __float2half(NAN);
|
||||
data[20] = __float2half(INFINITY);
|
||||
data[30] = __float2half(-INFINITY);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成包含NaN和Inf的测试数据 (float版本)
|
||||
std::vector<float> generateSpecialFloats(int size) {
|
||||
std::vector<float> data = generateRandomFloats(size, -10.0f, 10.0f);
|
||||
if (size > 100) {
|
||||
data[10] = NAN;
|
||||
data[20] = INFINITY;
|
||||
data[30] = -INFINITY;
|
||||
}
|
||||
return data;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 性能测试工具
|
||||
// ============================================================================
|
||||
class PerformanceMeter {
|
||||
private:
|
||||
mcEvent_t start, stop;
|
||||
|
||||
public:
|
||||
PerformanceMeter() {
|
||||
MACA_CHECK(mcEventCreate(&start));
|
||||
MACA_CHECK(mcEventCreate(&stop));
|
||||
}
|
||||
|
||||
~PerformanceMeter() {
|
||||
mcEventDestroy(start);
|
||||
mcEventDestroy(stop);
|
||||
}
|
||||
|
||||
void startTiming() {
|
||||
MACA_CHECK(mcEventRecord(start));
|
||||
}
|
||||
|
||||
float stopTiming() {
|
||||
MACA_CHECK(mcEventRecord(stop));
|
||||
MACA_CHECK(mcEventSynchronize(stop));
|
||||
float milliseconds = 0;
|
||||
MACA_CHECK(mcEventElapsedTime(&milliseconds, start, stop));
|
||||
return milliseconds;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 正确性验证工具
|
||||
// ============================================================================
|
||||
template<typename T>
|
||||
bool compareArrays(const std::vector<T>& a, const std::vector<T>& b, double tolerance = 1e-6) {
|
||||
if (a.size() != b.size()) return false;
|
||||
|
||||
for (size_t i = 0; i < a.size(); i++) {
|
||||
if constexpr (std::is_same_v<T, half>) {
|
||||
float fa = __half2float(a[i]);
|
||||
float fb = __half2float(b[i]);
|
||||
if (std::isnan(fa) && std::isnan(fb)) continue;
|
||||
if (std::isinf(fa) && std::isinf(fb) && (fa > 0) == (fb > 0)) continue;
|
||||
if (std::abs(fa - fb) > tolerance) return false;
|
||||
} else if constexpr (std::is_floating_point_v<T>) {
|
||||
if (std::isnan(a[i]) && std::isnan(b[i])) continue;
|
||||
if (std::isinf(a[i]) && std::isinf(b[i]) && (a[i] > 0) == (b[i] > 0)) continue;
|
||||
if (std::abs(a[i] - b[i]) > tolerance) return false;
|
||||
} else {
|
||||
if (a[i] != b[i]) return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
// CPU参考实现 - 稳定排序
|
||||
template<typename KeyType, typename ValueType>
|
||||
void cpuSortPair(std::vector<KeyType>& keys, std::vector<ValueType>& values, bool descending) {
|
||||
std::vector<std::pair<KeyType, ValueType>> pairs;
|
||||
for (size_t i = 0; i < keys.size(); i++) {
|
||||
pairs.emplace_back(keys[i], values[i]);
|
||||
}
|
||||
|
||||
if (descending) {
|
||||
std::stable_sort(pairs.begin(), pairs.end(),
|
||||
[](const auto& a, const auto& b) { return a.first > b.first; });
|
||||
} else {
|
||||
std::stable_sort(pairs.begin(), pairs.end());
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < pairs.size(); i++) {
|
||||
keys[i] = pairs[i].first;
|
||||
values[i] = pairs[i].second;
|
||||
}
|
||||
}
|
||||
|
||||
// CPU参考实现 - TopK
|
||||
template<typename KeyType, typename ValueType>
|
||||
void cpuTopkPair(const std::vector<KeyType>& keys_in, const std::vector<ValueType>& values_in,
|
||||
std::vector<KeyType>& keys_out, std::vector<ValueType>& values_out,
|
||||
int k, bool descending) {
|
||||
std::vector<std::pair<KeyType, ValueType>> pairs;
|
||||
for (size_t i = 0; i < keys_in.size(); i++) {
|
||||
pairs.emplace_back(keys_in[i], values_in[i]);
|
||||
}
|
||||
|
||||
if (descending) {
|
||||
std::stable_sort(pairs.begin(), pairs.end(),
|
||||
[](const auto& a, const auto& b) { return a.first > b.first; });
|
||||
} else {
|
||||
std::stable_sort(pairs.begin(), pairs.end());
|
||||
}
|
||||
|
||||
keys_out.resize(k);
|
||||
values_out.resize(k);
|
||||
for (int i = 0; i < k; i++) {
|
||||
keys_out[i] = pairs[i].first;
|
||||
values_out[i] = pairs[i].second;
|
||||
}
|
||||
}
|
||||
|
||||
// CPU参考实现 - ReduceSum (使用double精度)
|
||||
template<typename InputT>
|
||||
double cpuReduceSum(const std::vector<InputT>& data, double init_value) {
|
||||
double sum = init_value;
|
||||
for (const auto& val : data) {
|
||||
if constexpr (std::is_same_v<InputT, half>) {
|
||||
float f_val = __half2float(val);
|
||||
if (!std::isnan(f_val)) {
|
||||
sum += static_cast<double>(f_val);
|
||||
}
|
||||
} else {
|
||||
if (!std::isnan(val)) {
|
||||
sum += static_cast<double>(val);
|
||||
}
|
||||
}
|
||||
}
|
||||
return sum;
|
||||
}
|
||||
|
|
@ -0,0 +1,154 @@
|
|||
#pragma once
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <map>
|
||||
#include <string>
|
||||
#include <chrono>
|
||||
#include <iomanip>
|
||||
#include <sstream>
|
||||
|
||||
// ============================================================================
|
||||
// YAML性能报告生成器
|
||||
// ============================================================================
|
||||
|
||||
class YAMLPerformanceReporter {
|
||||
public:
|
||||
struct PerformanceData {
|
||||
std::string algorithm;
|
||||
std::string input_type;
|
||||
std::string output_type;
|
||||
std::string key_type;
|
||||
std::string value_type;
|
||||
std::vector<std::map<std::string, std::string>> metrics;
|
||||
};
|
||||
|
||||
// 创建性能数据条目
|
||||
static std::map<std::string, std::string> createEntry() {
|
||||
return std::map<std::string, std::string>();
|
||||
}
|
||||
|
||||
// 生成ReduceSum性能YAML
|
||||
static void generateReduceSumYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "reduce_sum_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "ReduceSum算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"ReduceSum\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " input: \"float\"\n";
|
||||
yaml_file << " output: \"float\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("throughput_gps")) << "\n";
|
||||
yaml_file << " data_type: \"" << data.at("data_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
// 生成SortPair性能YAML
|
||||
static void generateSortPairYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "sort_pair_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "SortPair算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"SortPair\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " key_type: \"float\"\n";
|
||||
yaml_file << " value_type: \"uint32_t\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " ascending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("asc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("asc_throughput_gps")) << "\n";
|
||||
yaml_file << " descending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("desc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("desc_throughput_gps")) << "\n";
|
||||
yaml_file << " key_type: \"" << data.at("key_type") << "\"\n";
|
||||
yaml_file << " value_type: \"" << data.at("value_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
// 生成TopkPair性能YAML
|
||||
static void generateTopkPairYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "topk_pair_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "TopkPair算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"TopkPair\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " key_type: \"float\"\n";
|
||||
yaml_file << " value_type: \"uint32_t\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " k_value: " << data.at("k_value") << "\n";
|
||||
yaml_file << " ascending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("asc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("asc_throughput_gps")) << "\n";
|
||||
yaml_file << " descending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("desc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("desc_throughput_gps")) << "\n";
|
||||
yaml_file << " key_type: \"" << data.at("key_type") << "\"\n";
|
||||
yaml_file << " value_type: \"" << data.at("value_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
private:
|
||||
// 写入YAML文件头部
|
||||
static void writeHeader(std::ofstream& file, const std::string& title) {
|
||||
file << "# " << title << "\n";
|
||||
file << "# 生成时间: ";
|
||||
|
||||
auto now = std::chrono::system_clock::now();
|
||||
auto time_t = std::chrono::system_clock::to_time_t(now);
|
||||
file << std::put_time(std::localtime(&time_t), "%Y-%m-%d %H:%M:%S");
|
||||
file << "\n\n";
|
||||
}
|
||||
|
||||
// 格式化浮点数
|
||||
static std::string formatFloat(const std::string& value) {
|
||||
try {
|
||||
double d = std::stod(value);
|
||||
std::ostringstream oss;
|
||||
oss << std::fixed << std::setprecision(6) << d;
|
||||
return oss.str();
|
||||
} catch (...) {
|
||||
return value;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
|
@ -0,0 +1,117 @@
|
|||
# GPU 高性能并行计算算法优化竞赛
|
||||
|
||||
## 📝 参赛指南
|
||||
|
||||
### 实现位置
|
||||
参赛者需要在以下文件中替换Thrust实现：
|
||||
- `reduce_sum_algorithm.maca` - 替换Thrust归约求和
|
||||
- `sort_pair_algorithm.maca` - 替换Thrust稳定排序
|
||||
- `topk_pair_algorithm.maca` - 替换Thrust TopK选择
|
||||
|
||||
### 算法要求
|
||||
见competition_parallel_algorithms.md
|
||||
|
||||
## 📊 性能评测
|
||||
|
||||
### 测试流程
|
||||
1. **Warmup**: 5次预热运行
|
||||
2. **Benchmark**: 10次正式测试取平均
|
||||
3. **数据规模**: 1M, 128M, 512M, 1G elements
|
||||
4. **评估指标**: 吞吐量(G/s)
|
||||
|
||||
### 性能指标计算
|
||||
|
||||
#### ReduceSum
|
||||
- **数据类型**: float → float
|
||||
- **吞吐量**: elements / time(s) / 1e9 (G/s)
|
||||
|
||||
#### SortPair
|
||||
- **数据类型**: <float, uint32_t>
|
||||
- **吞吐量**: elements / time(s) / 1e9 (G/s)
|
||||
|
||||
#### TopkPair
|
||||
- **数据类型**: <float, uint32_t>
|
||||
- **吞吐量**: elements / time(s) / 1e9 (G/s)
|
||||
|
||||
### 性能结果文件
|
||||
每个算法会生成详细的YAML性能分析文件：
|
||||
- `reduce_sum_performance.yaml` - ReduceSum性能数据
|
||||
- `sort_pair_performance.yaml` - SortPair性能数据
|
||||
- `topk_pair_performance.yaml` - TopkPair性能数据
|
||||
|
||||
这些文件包含：
|
||||
- 算法信息和数据类型
|
||||
- 计算公式说明
|
||||
- 各数据规模的详细性能数据
|
||||
- 升序/降序分别统计（适用时）
|
||||
|
||||
## 📁 提交内容结构
|
||||
|
||||
```
|
||||
├── run.sh # 统一编译和运行脚本（默认编译+运行所有算法）
|
||||
├── competition_parallel_algorithms.md # 详细题目说明
|
||||
│── reduce_sum_algorithm.maca # 1. ReduceSum测试程序
|
||||
│── sort_pair_algorithm.maca # 2. SortPair测试程序
|
||||
│── topk_pair_algorithm.maca # 3. TopkPair测试程序
|
||||
├── utils/ # 工具文件
|
||||
│ ├── test_utils.h # 测试工具和CPU参考实现
|
||||
│ ├── yaml_reporter.h # YAML性能报告生成器
|
||||
│ └── performance_utils.h # 性能测试工具
|
||||
├── reduce_sum_results.yaml #ReduceSum性能数据
|
||||
├── sort_pair_results.yaml #替换Thrust稳定排序
|
||||
└── topk_pair_results.yaml #TopkPair性能数据
|
||||
```
|
||||
|
||||
## 🔧 开发工具
|
||||
|
||||
### 编译选项
|
||||
```bash
|
||||
# 默认编译命令
|
||||
mxcc -O3 -std=c++17 --extended-lambda -Isrc
|
||||
|
||||
### 自动化测试
|
||||
```bash
|
||||
# 查看所有选项
|
||||
./build.sh --help
|
||||
|
||||
# 运行所有测试并生成YAML报告
|
||||
./build.sh --run_all
|
||||
|
||||
### 环境变量配置
|
||||
|
||||
| 变量 | 默认值 | 说明 |
|
||||
|--------|--------|------|
|
||||
| `COMPILER` | `mxcc` | CUDA编译器路径 |
|
||||
| `COMPILER_FLAGS` | `-O3 -std=c++17 --extended-lambda` | 编译标志 |
|
||||
| `HEADER_DIR` | `utils` | 头文件目录 |
|
||||
| `BUILD_DIR` | `build` | 构建输出目录 |
|
||||
|
||||
### 调试模式
|
||||
|
||||
## 📋 提交清单
|
||||
|
||||
在提交前请确保：
|
||||
- [ ] 所有算法通过正确性测试
|
||||
- [ ] 性能测试可以正常运行
|
||||
- [ ] 代码注释清晰，说明优化策略
|
||||
- [ ] 无内存泄漏或运行时错误
|
||||
- [ ] 生成完整测试报告
|
||||
- [ ] 在函数实现注释中说明创新点
|
||||
|
||||
# 提交时包含以下文件
|
||||
# - final_results/reduce_sum_results.yaml
|
||||
# - final_results/sort_pair_results.yaml
|
||||
# - final_results/topk_pair_results.yaml
|
||||
```
|
||||
|
||||
## 🤝 技术支持
|
||||
|
||||
如有技术问题，请：
|
||||
1. 查看详细错误信息和GPU状态
|
||||
2. 确认环境配置正确
|
||||
3. 检查内存使用是否超限
|
||||
4. 验证算法逻辑和数据类型
|
||||
|
||||
---
|
||||
|
||||
**祝您在竞赛中取得优异成绩！** 🏆
|
||||
|
|
@ -0,0 +1,274 @@
|
|||
#!/bin/bash
|
||||
|
||||
# GPU高性能并行计算算法优化竞赛 - 统一编译和运行脚本
|
||||
# 整合了所有算法的编译、运行和公共配置
|
||||
|
||||
# ============================================================================
|
||||
# 公共配置和工具函数
|
||||
# ============================================================================
|
||||
|
||||
# 设置颜色
|
||||
RED='\033[0;31m'
|
||||
GREEN='\033[0;32m'
|
||||
BLUE='\033[0;34m'
|
||||
YELLOW='\033[0;33m'
|
||||
NC='\033[0m' # No Color
|
||||
|
||||
# 打印函数
|
||||
print_info() {
|
||||
echo -e "${BLUE}[INFO]${NC} $1"
|
||||
}
|
||||
|
||||
print_success() {
|
||||
echo -e "${GREEN}[SUCCESS]${NC} $1"
|
||||
}
|
||||
|
||||
print_error() {
|
||||
echo -e "${RED}[ERROR]${NC} $1"
|
||||
}
|
||||
|
||||
print_warning() {
|
||||
echo -e "${YELLOW}[WARNING]${NC} $1"
|
||||
}
|
||||
|
||||
# 编译配置 - 可通过环境变量自定义
|
||||
COMPILER=${COMPILER:-mxcc}
|
||||
COMPILER_FLAGS=${COMPILER_FLAGS:-"-O3 -std=c++17 --extended-lambda -DRUN_FULL_TEST"}
|
||||
|
||||
# ***** 这里是关键修改点1：头文件目录 *****
|
||||
# 现在头文件在 utils/ 目录下
|
||||
HEADER_DIR=${HEADER_DIR:-utils}
|
||||
|
||||
# ***** 这里是关键修改点2：源文件目录 *****
|
||||
# 现在源文件在 ./ 目录下
|
||||
SOURCE_CODE_DIR=${SOURCE_CODE_DIR:-}
|
||||
|
||||
BUILD_DIR=${BUILD_DIR:-build}
|
||||
|
||||
# 编译单个算法的通用函数
|
||||
# 参数: $1=算法名称, $2=源文件名（不含路径）
|
||||
compile_algorithm() {
|
||||
local algo_name="$1"
|
||||
local source_file_name="$2" # 例如 "reduce_sum_algorithm.maca"
|
||||
local target_file="$BUILD_DIR/test_${algo_name,,}" # 转换为小写
|
||||
|
||||
print_info "编译 $algo_name 算法..."
|
||||
|
||||
# 创建构建目录
|
||||
mkdir -p "$BUILD_DIR"
|
||||
|
||||
# ***** 这里是关键修改点3：编译命令 *****
|
||||
# -I$HEADER_DIR 用于告诉编译器头文件在哪里
|
||||
# $SOURCE_CODE_DIR/$source_file_name 用于指定要编译的源文件的完整路径
|
||||
local compile_cmd="$COMPILER $COMPILER_FLAGS -I$HEADER_DIR $source_file_name -o $target_file"
|
||||
|
||||
print_info "执行: $compile_cmd"
|
||||
|
||||
if $compile_cmd; then
|
||||
print_success "$algo_name 编译完成!"
|
||||
echo ""
|
||||
echo "运行测试:"
|
||||
echo " ./$target_file [correctness|performance|all]"
|
||||
return 0
|
||||
else
|
||||
print_error "$algo_name 编译失败!"
|
||||
return 1
|
||||
fi
|
||||
}
|
||||
|
||||
# 显示编译配置信息
|
||||
show_build_config() {
|
||||
print_info "编译配置:"
|
||||
echo " COMPILER: $COMPILER"
|
||||
echo " COMPILER_FLAGS: $COMPILER_FLAGS"
|
||||
echo " HEADER_DIR: $HEADER_DIR" # 显示头文件目录
|
||||
echo " SOURCE_CODE_DIR: $SOURCE_CODE_DIR" # 显示源文件目录
|
||||
echo " BUILD_DIR: $BUILD_DIR"
|
||||
echo ""
|
||||
}
|
||||
|
||||
# 运行单个测试
|
||||
run_single_test() {
|
||||
local algo_name="$1"
|
||||
local test_mode="${2:-all}"
|
||||
local test_file="$BUILD_DIR/test_${algo_name,,}"
|
||||
|
||||
if [ -f "$test_file" ]; then
|
||||
print_info "运行 $algo_name 测试 (模式: $test_mode)..."
|
||||
"./$test_file" "$test_mode"
|
||||
return $?
|
||||
else
|
||||
print_error "$algo_name 测试程序不存在: $test_file"
|
||||
return 1
|
||||
fi
|
||||
}
|
||||
|
||||
# ============================================================================
|
||||
# 主脚本逻辑
|
||||
# ============================================================================
|
||||
|
||||
# 显示帮助信息 (整合了所有选项)
|
||||
show_help() {
|
||||
echo "GPU算法竞赛统一编译和运行脚本"
|
||||
echo "用法: $0 [选项]"
|
||||
echo ""
|
||||
echo "选项:"
|
||||
echo " --help 显示帮助信息"
|
||||
echo " --build-only 仅编译所有算法，不运行测试"
|
||||
echo " --run_reduce [MODE] 编译并运行ReduceSum算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo " --run_sort [MODE] 编译并运行SortPair算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo " --run_topk [MODE] 编译并运行TopkPair算法测试 (MODE: correctness|performance|all, 默认all)"
|
||||
echo ""
|
||||
echo "示例:"
|
||||
echo " $0 # 编译并运行所有测试（默认行为）"
|
||||
echo " $0 --build-only # 仅编译所有算法"
|
||||
echo " $0 --run_sort performance # 编译并运行SortPair性能测试"
|
||||
echo ""
|
||||
}
|
||||
|
||||
# 解析命令行参数
|
||||
RUN_MODE="run_all" # 默认为编译并运行所有测试
|
||||
ALGO_TO_RUN="" # 记录要运行的单个算法
|
||||
SINGLE_ALGO_TEST_MODE="all" # 单个算法的测试模式
|
||||
|
||||
while [[ $# -gt 0 ]]; do
|
||||
case $1 in
|
||||
--help)
|
||||
show_help
|
||||
exit 0
|
||||
;;
|
||||
--build-only)
|
||||
RUN_MODE="build_only"
|
||||
shift
|
||||
;;
|
||||
--run_reduce)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="ReduceSum"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
--run_sort)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="SortPair"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
--run_topk)
|
||||
RUN_MODE="run_single"
|
||||
ALGO_TO_RUN="TopkPair"
|
||||
if [[ -n "$2" && "$2" != --* ]]; then
|
||||
SINGLE_ALGO_TEST_MODE="$2"
|
||||
shift
|
||||
fi
|
||||
shift
|
||||
;;
|
||||
*)
|
||||
print_error "未知选项: $1"
|
||||
show_help
|
||||
exit 1
|
||||
;;
|
||||
esac
|
||||
done
|
||||
|
||||
if [ "$RUN_MODE" = "build_only" ]; then
|
||||
print_info "开始编译所有算法..."
|
||||
else
|
||||
print_info "开始编译并运行所有算法..."
|
||||
fi
|
||||
print_info "工作目录: $(pwd)"
|
||||
print_info "编译时间: $(date '+%Y-%m-%d %H:%M:%S')"
|
||||
show_build_config
|
||||
|
||||
# 清理构建目录
|
||||
if [ -d "$BUILD_DIR" ]; then
|
||||
print_info "清理现有构建目录: $BUILD_DIR"
|
||||
rm -rf "$BUILD_DIR"
|
||||
fi
|
||||
|
||||
# 核心逻辑：根据 RUN_MODE 执行操作
|
||||
case "$RUN_MODE" in
|
||||
"build_only")
|
||||
print_info "编译所有算法..."
|
||||
|
||||
# 直接调用 compile_algorithm 函数
|
||||
print_info "[1/3] 编译ReduceSum..."
|
||||
if ! compile_algorithm "ReduceSum" "reduce_sum_algorithm.maca"; then
|
||||
print_error "ReduceSum编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[2/3] 编译SortPair..."
|
||||
if ! compile_algorithm "SortPair" "sort_pair_algorithm.maca"; then
|
||||
print_error "SortPair编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[3/3] 编译TopkPair..."
|
||||
if ! compile_algorithm "TopkPair" "topk_pair_algorithm.maca"; then
|
||||
print_error "TopkPair编译失败"
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_success "所有算法编译完成!"
|
||||
echo ""
|
||||
echo "可执行文件:"
|
||||
echo " $BUILD_DIR/test_reducesum - ReduceSum算法测试"
|
||||
echo " $BUILD_DIR/test_sortpair - SortPair算法测试"
|
||||
echo " $BUILD_DIR/test_topkpair - TopkPair算法测试"
|
||||
echo ""
|
||||
echo "使用方法:"
|
||||
echo " ./$BUILD_DIR/test_reducesum [correctness|performance|all]"
|
||||
echo " ./$BUILD_DIR/test_sortpair [correctness|performance|all]"
|
||||
echo " ./$BUILD_DIR/test_topkpair [correctness|performance|all]"
|
||||
;;
|
||||
|
||||
"run_all")
|
||||
print_info "编译并运行所有算法测试..."
|
||||
|
||||
# 直接调用 compile_algorithm 和 run_single_test 函数
|
||||
print_info "[1/3] ReduceSum..."
|
||||
if compile_algorithm "ReduceSum" "reduce_sum_algorithm.maca"; then
|
||||
run_single_test "ReduceSum" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[2/3] SortPair..."
|
||||
if compile_algorithm "SortPair" "sort_pair_algorithm.maca"; then
|
||||
run_single_test "SortPair" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_info "[3/3] TopkPair..."
|
||||
if compile_algorithm "TopkPair" "topk_pair_algorithm.maca"; then
|
||||
run_single_test "TopkPair" "all"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
|
||||
print_success "所有测试完成!"
|
||||
;;
|
||||
|
||||
"run_single")
|
||||
print_info "编译并运行 ${ALGO_TO_RUN} 测试 (模式: ${SINGLE_ALGO_TEST_MODE})..."
|
||||
local source_file_name=""
|
||||
case "$ALGO_TO_RUN" in
|
||||
"ReduceSum") source_file_name="reduce_sum_algorithm.maca" ;;
|
||||
"SortPair") source_file_name="sort_pair_algorithm.maca" ;;
|
||||
"TopkPair") source_file_name="topk_pair_algorithm.maca" ;;
|
||||
esac
|
||||
|
||||
if compile_algorithm "$ALGO_TO_RUN" "$source_file_name"; then
|
||||
run_single_test "$ALGO_TO_RUN" "$SINGLE_ALGO_TEST_MODE"
|
||||
else
|
||||
exit 1
|
||||
fi
|
||||
;;
|
||||
esac
|
||||
|
|
@ -0,0 +1,97 @@
|
|||
# 样例赛题说明
|
||||
|
||||
## GPU高性能并行计算算法优化
|
||||
|
||||
要求参赛者通过一个或多个global kernel 函数（允许配套 device 辅助函数），实现高性能算法。
|
||||
|
||||
在正确性、稳定性前提下，比拼算法性能。
|
||||
|
||||
# 1. ReduceSum算法优化
|
||||
```cpp
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
class ReduceSumAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：InputT: float, OutputT: float
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入d\_in数据量为num\_items
|
||||
|
||||
注意事项
|
||||
|
||||
* 累计误差不大于cpu double golden基准的0.5%
|
||||
* 注意针对NAN和INF等异常值的处理
|
||||
|
||||
|
||||
加分项
|
||||
|
||||
* 使用tensor core计算reduce
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
||||
# 2. Sort Pair算法优化
|
||||
```cpp
|
||||
template <typename KeyType, typename ValueType>
|
||||
class SortPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void sort(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, bool descending) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：key: float, value: int32\_t
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入、输出的key和value的数据量一致，均为num\_items
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 需要校验结果正确性
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的排序，如half、double、int32_t等
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
# 3. Topk Pair算法优化
|
||||
```cpp
|
||||
template <typename KeyType, typename ValueType>
|
||||
class TopkPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void topk(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, int k, bool descending) {
|
||||
// TODO
|
||||
}
|
||||
};
|
||||
```
|
||||
其中
|
||||
|
||||
* 数据类型：key: float, value: int32\_t
|
||||
* 系统将测试评估1M, 128M, 512M, 1G element number下的算法性能
|
||||
* 假定输入的key和value的数据量一致，为num\_items；输出的key和value的数据量一致，为k
|
||||
* k的范围：32，50，100，256，1024。k不大于num\_items
|
||||
|
||||
|
||||
注意事项
|
||||
|
||||
* 结果必须稳定排序
|
||||
|
||||
加分项
|
||||
|
||||
* 支持其他不同数据类型的键值对，实现类型通用算法
|
||||
* 覆盖更全面的数据范围，提供良好稳定的性能表现
|
||||
|
||||
|
|
@ -0,0 +1,277 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/reduce.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/functional.h>
|
||||
#endif
|
||||
|
||||
// 误差容忍度
|
||||
constexpr double REDUCE_ERROR_TOLERANCE = 0.005; // 0.5%
|
||||
|
||||
// ============================================================================
|
||||
// ReduceSum算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename InputT = float, typename OutputT = float>
|
||||
class ReduceSumAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void reduce(const InputT* d_in, OutputT* d_out, int num_items, OutputT init_value) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能归约算法
|
||||
|
||||
// 示例：参赛者可以调用1个或多个自定义kernel
|
||||
// blockReduceKernel<<<grid, block>>>(d_in, temp_results, num_items, init_value);
|
||||
// finalReduceKernel<<<1, block>>>(temp_results, d_out, grid.x);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
auto input_ptr = thrust::device_pointer_cast(d_in);
|
||||
auto output_ptr = thrust::device_pointer_cast(d_out);
|
||||
|
||||
// 直接使用thrust::reduce进行归约
|
||||
*output_ptr = thrust::reduce(
|
||||
thrust::device,
|
||||
input_ptr,
|
||||
input_ptr + num_items,
|
||||
static_cast<OutputT>(init_value)
|
||||
);
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：中间结果缓冲区、多阶段归约等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "ReduceSum 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
ReduceSumAlgorithm<float, float> algorithm;
|
||||
|
||||
bool allPassed = true;
|
||||
|
||||
// 测试不同数据规模
|
||||
for (int i = 0; i < NUM_TEST_SIZES && i < 2; i++) { // 限制测试规模
|
||||
int size = std::min(TEST_SIZES[i], 10000);
|
||||
std::cout << " 测试规模: " << size << std::endl;
|
||||
|
||||
// 测试普通数据
|
||||
{
|
||||
auto data = generator.generateRandomFloats(size, -10.0f, 10.0f);
|
||||
float init_value = 1.0f;
|
||||
|
||||
// CPU参考计算
|
||||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||||
|
||||
// GPU计算
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证误差
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||||
std::cout << " 失败: 误差过大 " << relative_error << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过 (误差: " << relative_error << ")" << std::endl;
|
||||
}
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
|
||||
// 测试特殊值 (NaN, Inf)
|
||||
if (size > 100) {
|
||||
std::cout << " 测试特殊值..." << std::endl;
|
||||
auto data = generator.generateSpecialFloats(size);
|
||||
float init_value = 0.0f;
|
||||
|
||||
double cpu_result = cpuReduceSum(data, static_cast<double>(init_value));
|
||||
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
|
||||
float gpu_result;
|
||||
MACA_CHECK(mcMemcpy(&gpu_result, d_out, sizeof(float), mcMemcpyDeviceToHost));
|
||||
|
||||
// 对于包含特殊值的情况，检查是否正确处理
|
||||
if (std::isfinite(cpu_result) && std::isfinite(gpu_result)) {
|
||||
double relative_error = std::abs(gpu_result - cpu_result) / std::abs(cpu_result);
|
||||
if (relative_error > REDUCE_ERROR_TOLERANCE) {
|
||||
std::cout << " 失败: 特殊值处理错误" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过 (特殊值处理)" << std::endl;
|
||||
}
|
||||
} else {
|
||||
std::cout << " 通过 (特殊值结果)" << std::endl;
|
||||
}
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
}
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
PerformanceDisplay::printReduceSumHeader();
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
ReduceSumAlgorithm<float, float> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
// 生成测试数据
|
||||
auto data = generator.generateRandomFloats(size);
|
||||
float init_value = 0.0f;
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_in;
|
||||
float *d_out;
|
||||
MACA_CHECK(mcMalloc(&d_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_out, sizeof(float)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_in, data.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.reduce(d_in, d_out, size, init_value);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
|
||||
// 计算性能指标
|
||||
auto metrics = PerformanceCalculator::calculateReduceSum(size, avg_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printReduceSumData(size, avg_time, metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["time_ms"] = std::to_string(avg_time);
|
||||
entry["throughput_gps"] = std::to_string(metrics.throughput_gps);
|
||||
entry["data_type"] = "float";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
mcFree(d_in);
|
||||
mcFree(d_out);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateReduceSumYAML(perf_data, "reduce_sum_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("reduce_sum_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== ReduceSum 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << ReduceSumAlgorithm<float, float>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,13 @@
|
|||
#!/bin/bash
|
||||
|
||||
# 单个赛题测试验证(ReduceSum算法)
|
||||
#./build_and_run.sh --run_reduce
|
||||
|
||||
# 单个赛题测试验证(SortPair算法)
|
||||
#./build_and_run.sh --run_reduce
|
||||
|
||||
# 单个赛题测试验证(TopkPair算法)
|
||||
# ./build_and_run.sh --run_topk
|
||||
|
||||
# 默认全量赛题测试验证，参赛选手单个优化，参考单个脚本执行方式，CI入口run.sh
|
||||
./build_and_run.sh
|
||||
|
|
@ -0,0 +1,275 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/iterator/zip_iterator.h>
|
||||
#include <thrust/tuple.h>
|
||||
#endif
|
||||
|
||||
// ============================================================================
|
||||
// SortPair算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename KeyType, typename ValueType>
|
||||
class SortPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void sort(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, bool descending) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能排序算法
|
||||
|
||||
// 示例：参赛者可以调用1个或多个自定义kernel
|
||||
// preprocessKernel<<<grid, block>>>(d_keys_in, d_values_in, num_items);
|
||||
// mainSortKernel<<<grid, block>>>(d_keys_out, d_values_out, num_items, descending);
|
||||
// postprocessKernel<<<grid, block>>>(d_keys_out, d_values_out, num_items);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_out, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_out, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
auto key_ptr = thrust::device_pointer_cast(d_keys_out);
|
||||
auto value_ptr = thrust::device_pointer_cast(d_values_out);
|
||||
|
||||
if (descending) {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：临时缓冲区、多个kernel函数、流等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "SortPair 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
SortPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
// 测试小规模数据
|
||||
int size = 10000;
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
// 测试升序和降序
|
||||
bool allPassed = true;
|
||||
for (bool descending : {false, true}) {
|
||||
std::cout << " " << (descending ? "降序" : "升序") << " 测试..." << std::endl;
|
||||
|
||||
// CPU参考结果
|
||||
auto cpu_keys = keys;
|
||||
auto cpu_values = values;
|
||||
cpuSortPair(cpu_keys, cpu_values, descending);
|
||||
|
||||
// GPU算法结果
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
|
||||
// 获取结果
|
||||
std::vector<float> gpu_keys(size);
|
||||
std::vector<uint32_t> gpu_values(size);
|
||||
MACA_CHECK(mcMemcpy(gpu_keys.data(), d_keys_out, size * sizeof(float), mcMemcpyDeviceToHost));
|
||||
MACA_CHECK(mcMemcpy(gpu_values.data(), d_values_out, size * sizeof(uint32_t), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证结果
|
||||
bool keysMatch = compareArrays(cpu_keys, gpu_keys, 1e-5);
|
||||
bool valuesMatch = compareArrays(cpu_values, gpu_values);
|
||||
|
||||
if (!keysMatch || !valuesMatch) {
|
||||
std::cout << " 失败: 结果不匹配" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_in);
|
||||
mcFree(d_values_out);
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
PerformanceDisplay::printSortPairHeader();
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
SortPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
for (int i = 0; i < NUM_TEST_SIZES; i++) {
|
||||
int size = TEST_SIZES[i];
|
||||
|
||||
// 生成测试数据
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
float asc_time = 0, desc_time = 0;
|
||||
|
||||
// 测试升序和降序
|
||||
for (bool descending : {false, true}) {
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.sort(d_keys_in, d_keys_out, d_values_in, d_values_out, size, descending);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
if (descending) {
|
||||
desc_time = avg_time;
|
||||
} else {
|
||||
asc_time = avg_time;
|
||||
}
|
||||
}
|
||||
|
||||
// 计算性能指标
|
||||
auto asc_metrics = PerformanceCalculator::calculateSortPair(size, asc_time);
|
||||
auto desc_metrics = PerformanceCalculator::calculateSortPair(size, desc_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printSortPairData(size, asc_time, desc_time, asc_metrics, desc_metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["asc_time_ms"] = std::to_string(asc_time);
|
||||
entry["desc_time_ms"] = std::to_string(desc_time);
|
||||
entry["asc_throughput_gps"] = std::to_string(asc_metrics.throughput_gps);
|
||||
entry["desc_throughput_gps"] = std::to_string(desc_metrics.throughput_gps);
|
||||
entry["key_type"] = "float";
|
||||
entry["value_type"] = "uint32_t";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_in);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateSortPairYAML(perf_data, "sort_pair_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("sort_pair_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== SortPair 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << SortPairAlgorithm<float, uint32_t>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,317 @@
|
|||
#include "test_utils.h"
|
||||
#include "performance_utils.h"
|
||||
#include "yaml_reporter.h"
|
||||
#include <iostream>
|
||||
#include <vector>
|
||||
#include <iomanip>
|
||||
#include <fstream>
|
||||
#include <map>
|
||||
#include <chrono>
|
||||
|
||||
// ============================================================================
|
||||
// 实现标记宏 - 参赛者修改实现时请将此宏设为0
|
||||
// ============================================================================
|
||||
#ifndef USE_DEFAULT_REF_IMPL
|
||||
#define USE_DEFAULT_REF_IMPL 1 // 1=默认实现, 0=参赛者自定义实现
|
||||
#endif
|
||||
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
#include <thrust/sort.h>
|
||||
#include <thrust/device_vector.h>
|
||||
#include <thrust/execution_policy.h>
|
||||
#include <thrust/iterator/zip_iterator.h>
|
||||
#include <thrust/tuple.h>
|
||||
#include <thrust/copy.h>
|
||||
#endif
|
||||
|
||||
static const int TOPK_VALUES[] = {32, 50, 100, 256, 1024};
|
||||
static const int NUM_TOPK_VALUES = sizeof(TOPK_VALUES) / sizeof(TOPK_VALUES[0]);
|
||||
|
||||
// ============================================================================
|
||||
// TopkPair算法实现接口
|
||||
// 参赛者需要替换Thrust实现为自己的高性能kernel
|
||||
// ============================================================================
|
||||
|
||||
template <typename KeyType, typename ValueType>
|
||||
class TopkPairAlgorithm {
|
||||
public:
|
||||
// 主要接口函数 - 参赛者需要实现这个函数
|
||||
void topk(const KeyType* d_keys_in, KeyType* d_keys_out,
|
||||
const ValueType* d_values_in, ValueType* d_values_out,
|
||||
int num_items, int k, bool descending) {
|
||||
|
||||
#if !USE_DEFAULT_REF_IMPL
|
||||
// ========================================
|
||||
// 参赛者自定义实现区域
|
||||
// ========================================
|
||||
|
||||
// TODO: 参赛者在此实现自己的高性能TopK算法
|
||||
|
||||
// 示例：参赛者可以调用多个自定义kernel
|
||||
// TopkKernel1<<<grid, block>>>(d_keys_in, d_values_in, temp_results, num_items, k);
|
||||
// TopkKernel2<<<grid, block>>>(temp_results, d_keys_out, d_values_out, k, descending);
|
||||
#else
|
||||
// ========================================
|
||||
// 默认基准实现
|
||||
// ========================================
|
||||
|
||||
KeyType* temp_keys;
|
||||
ValueType* temp_values;
|
||||
MACA_CHECK(mcMalloc(&temp_keys, num_items * sizeof(KeyType)));
|
||||
MACA_CHECK(mcMalloc(&temp_values, num_items * sizeof(ValueType)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(temp_keys, d_keys_in, num_items * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(temp_values, d_values_in, num_items * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
auto key_ptr = thrust::device_pointer_cast(temp_keys);
|
||||
auto value_ptr = thrust::device_pointer_cast(temp_values);
|
||||
|
||||
// 由于greater和less是不同类型，需要分别调用
|
||||
if (descending) {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::greater<KeyType>());
|
||||
} else {
|
||||
thrust::stable_sort_by_key(thrust::device, key_ptr, key_ptr + num_items, value_ptr, thrust::less<KeyType>());
|
||||
}
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_out, temp_keys, k * sizeof(KeyType), mcMemcpyDeviceToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_out, temp_values, k * sizeof(ValueType), mcMemcpyDeviceToDevice));
|
||||
|
||||
mcFree(temp_keys);
|
||||
mcFree(temp_values);
|
||||
#endif
|
||||
}
|
||||
|
||||
// 获取当前实现状态
|
||||
static const char* getImplementationStatus() {
|
||||
#if USE_DEFAULT_REF_IMPL
|
||||
return "DEFAULT_REF_IMPL";
|
||||
#else
|
||||
return "CUSTOM_IMPL";
|
||||
#endif
|
||||
}
|
||||
|
||||
private:
|
||||
// 参赛者可以在这里添加辅助函数和成员变量
|
||||
// 例如：分块大小、临时缓冲区、多流处理等
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 测试和性能评估
|
||||
// ============================================================================
|
||||
|
||||
bool testCorrectness() {
|
||||
std::cout << "TopkPair 正确性测试..." << std::endl;
|
||||
TestDataGenerator generator;
|
||||
TopkPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
int size = 10000;
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in, *d_keys_out;
|
||||
uint32_t *d_values_in, *d_values_out;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
bool allPassed = true;
|
||||
|
||||
// 测试不同k值
|
||||
for (int ki = 0; ki < NUM_TOPK_VALUES && ki < 4; ki++) { // 限制测试范围
|
||||
int k = TOPK_VALUES[ki];
|
||||
if (k > size) continue;
|
||||
|
||||
std::cout << " 测试 k=" << k << std::endl;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t)));
|
||||
|
||||
for (bool descending : {false, true}) {
|
||||
std::cout << " " << (descending ? "降序" : "升序") << " TopK..." << std::endl;
|
||||
|
||||
// CPU参考结果
|
||||
std::vector<float> cpu_keys_out;
|
||||
std::vector<uint32_t> cpu_values_out;
|
||||
cpuTopkPair(keys, values, cpu_keys_out, cpu_values_out, k, descending);
|
||||
|
||||
// GPU算法结果
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
|
||||
// 获取结果
|
||||
std::vector<float> gpu_keys_out(k);
|
||||
std::vector<uint32_t> gpu_values_out(k);
|
||||
MACA_CHECK(mcMemcpy(gpu_keys_out.data(), d_keys_out, k * sizeof(float), mcMemcpyDeviceToHost));
|
||||
MACA_CHECK(mcMemcpy(gpu_values_out.data(), d_values_out, k * sizeof(uint32_t), mcMemcpyDeviceToHost));
|
||||
|
||||
// 验证结果
|
||||
bool keysMatch = compareArrays(cpu_keys_out, gpu_keys_out, 1e-5);
|
||||
bool valuesMatch = compareArrays(cpu_values_out, gpu_values_out);
|
||||
|
||||
if (!keysMatch || !valuesMatch) {
|
||||
std::cout << " 失败: 结果不匹配" << std::endl;
|
||||
allPassed = false;
|
||||
} else {
|
||||
std::cout << " 通过" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
// 清理内存
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_values_in);
|
||||
|
||||
return allPassed;
|
||||
}
|
||||
|
||||
void benchmarkPerformance() {
|
||||
std::cout << "\nTopkPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
|
||||
TestDataGenerator generator;
|
||||
PerformanceMeter meter;
|
||||
TopkPairAlgorithm<float, uint32_t> algorithm;
|
||||
|
||||
const int WARMUP_ITERATIONS = 5;
|
||||
const int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
// 用于YAML报告的数据收集
|
||||
std::vector<std::map<std::string, std::string>> perf_data;
|
||||
|
||||
// 针对不同数据规模测试
|
||||
for (int size_idx = 0; size_idx < NUM_TEST_SIZES; size_idx++) {
|
||||
int size = TEST_SIZES[size_idx];
|
||||
std::cout << "\n数据规模: " << size << std::endl;
|
||||
std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(74, '-') << std::endl;
|
||||
|
||||
auto keys = generator.generateRandomFloats(size);
|
||||
auto values = generator.generateRandomUint32(size);
|
||||
|
||||
// 分配GPU内存
|
||||
float *d_keys_in;
|
||||
uint32_t *d_values_in;
|
||||
|
||||
MACA_CHECK(mcMalloc(&d_keys_in, size * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_in, size * sizeof(uint32_t)));
|
||||
|
||||
MACA_CHECK(mcMemcpy(d_keys_in, keys.data(), size * sizeof(float), mcMemcpyHostToDevice));
|
||||
MACA_CHECK(mcMemcpy(d_values_in, values.data(), size * sizeof(uint32_t), mcMemcpyHostToDevice));
|
||||
|
||||
for (int ki = 0; ki < NUM_TOPK_VALUES; ki++) {
|
||||
int k = TOPK_VALUES[ki];
|
||||
if (k > size) continue;
|
||||
|
||||
float *d_keys_out;
|
||||
uint32_t *d_values_out;
|
||||
MACA_CHECK(mcMalloc(&d_keys_out, k * sizeof(float)));
|
||||
MACA_CHECK(mcMalloc(&d_values_out, k * sizeof(uint32_t)));
|
||||
|
||||
float asc_time = 0, desc_time = 0;
|
||||
|
||||
for (bool descending : {false, true}) {
|
||||
// Warmup阶段
|
||||
for (int iter = 0; iter < WARMUP_ITERATIONS; iter++) {
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
}
|
||||
|
||||
// 正式测试阶段
|
||||
float total_time = 0;
|
||||
for (int iter = 0; iter < BENCHMARK_ITERATIONS; iter++) {
|
||||
meter.startTiming();
|
||||
algorithm.topk(d_keys_in, d_keys_out, d_values_in, d_values_out, size, k, descending);
|
||||
total_time += meter.stopTiming();
|
||||
}
|
||||
|
||||
float avg_time = total_time / BENCHMARK_ITERATIONS;
|
||||
if (descending) {
|
||||
desc_time = avg_time;
|
||||
} else {
|
||||
asc_time = avg_time;
|
||||
}
|
||||
}
|
||||
|
||||
// 计算性能指标
|
||||
auto asc_metrics = PerformanceCalculator::calculateTopkPair(size, k, asc_time);
|
||||
auto desc_metrics = PerformanceCalculator::calculateTopkPair(size, k, desc_time);
|
||||
|
||||
// 显示性能数据
|
||||
PerformanceDisplay::printTopkPairData(k, asc_time, desc_time, asc_metrics, desc_metrics);
|
||||
|
||||
// 收集YAML报告数据
|
||||
auto entry = YAMLPerformanceReporter::createEntry();
|
||||
entry["data_size"] = std::to_string(size);
|
||||
entry["k_value"] = std::to_string(k);
|
||||
entry["asc_time_ms"] = std::to_string(asc_time);
|
||||
entry["desc_time_ms"] = std::to_string(desc_time);
|
||||
entry["asc_throughput_gps"] = std::to_string(asc_metrics.throughput_gps);
|
||||
entry["desc_throughput_gps"] = std::to_string(desc_metrics.throughput_gps);
|
||||
entry["key_type"] = "float";
|
||||
entry["value_type"] = "uint32_t";
|
||||
perf_data.push_back(entry);
|
||||
|
||||
mcFree(d_keys_out);
|
||||
mcFree(d_values_out);
|
||||
}
|
||||
|
||||
mcFree(d_keys_in);
|
||||
mcFree(d_values_in);
|
||||
}
|
||||
|
||||
// 生成YAML性能报告
|
||||
YAMLPerformanceReporter::generateTopkPairYAML(perf_data, "topk_pair_performance.yaml");
|
||||
PerformanceDisplay::printSavedMessage("topk_pair_performance.yaml");
|
||||
}
|
||||
|
||||
// ============================================================================
|
||||
// 主函数
|
||||
// ============================================================================
|
||||
int main(int argc, char* argv[]) {
|
||||
std::cout << "=== TopkPair 算法测试 ===" << std::endl;
|
||||
|
||||
// 检查参数
|
||||
std::string mode = "all";
|
||||
if (argc > 1) {
|
||||
mode = argv[1];
|
||||
}
|
||||
|
||||
bool correctness_passed = true;
|
||||
bool performance_completed = true;
|
||||
|
||||
try {
|
||||
if (mode == "correctness" || mode == "all") {
|
||||
correctness_passed = testCorrectness();
|
||||
}
|
||||
|
||||
if (mode == "performance" || mode == "all") {
|
||||
if (correctness_passed || mode == "performance") {
|
||||
benchmarkPerformance();
|
||||
} else {
|
||||
std::cout << "跳过性能测试，因为正确性测试未通过" << std::endl;
|
||||
performance_completed = false;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "\n=== 测试完成 ===" << std::endl;
|
||||
std::cout << "实现状态: " << TopkPairAlgorithm<float, uint32_t>::getImplementationStatus() << std::endl;
|
||||
if (mode == "all") {
|
||||
std::cout << "正确性: " << (correctness_passed ? "通过" : "失败") << std::endl;
|
||||
std::cout << "性能测试: " << (performance_completed ? "完成" : "跳过") << std::endl;
|
||||
}
|
||||
|
||||
return correctness_passed ? 0 : 1;
|
||||
|
||||
} catch (const std::exception& e) {
|
||||
std::cerr << "测试出错: " << e.what() << std::endl;
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
|
|
@ -0,0 +1,114 @@
|
|||
#pragma once
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#include <string>
|
||||
|
||||
// ============================================================================
|
||||
// 性能计算和显示工具
|
||||
// ============================================================================
|
||||
|
||||
class PerformanceCalculator {
|
||||
public:
|
||||
// ReduceSum性能计算
|
||||
struct ReduceSumMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static ReduceSumMetrics calculateReduceSum(int size, float time_ms) {
|
||||
ReduceSumMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
|
||||
// SortPair性能计算
|
||||
struct SortPairMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static SortPairMetrics calculateSortPair(int size, float time_ms) {
|
||||
SortPairMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
|
||||
// TopkPair性能计算
|
||||
struct TopkPairMetrics {
|
||||
double throughput_gps; // G elements/s
|
||||
};
|
||||
|
||||
static TopkPairMetrics calculateTopkPair(int size, int k, float time_ms) {
|
||||
TopkPairMetrics metrics;
|
||||
metrics.throughput_gps = (size / 1e9) / (time_ms / 1000.0);
|
||||
return metrics;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 性能显示工具
|
||||
// ============================================================================
|
||||
|
||||
class PerformanceDisplay {
|
||||
public:
|
||||
// 显示ReduceSum性能表头
|
||||
static void printReduceSumHeader() {
|
||||
std::cout << "\nReduceSum 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: float -> float" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
std::cout << std::setw(12) << "数据规模" << std::setw(15) << "时间(ms)"
|
||||
<< std::setw(20) << "吞吐量(G/s)" << std::endl;
|
||||
std::cout << std::string(47, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示SortPair性能表头
|
||||
static void printSortPairHeader() {
|
||||
std::cout << "\nSortPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
std::cout << std::setw(12) << "数据规模" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(78, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示TopkPair性能表头
|
||||
static void printTopkPairHeader() {
|
||||
std::cout << "\nTopkPair 性能测试..." << std::endl;
|
||||
std::cout << "数据类型: <float, uint32_t>" << std::endl;
|
||||
std::cout << "计算公式:" << std::endl;
|
||||
std::cout << " 吞吐量 = 元素数 / 时间(s) / 1e9 (G/s)" << std::endl;
|
||||
}
|
||||
|
||||
static void printTopkPairDataHeader() {
|
||||
std::cout << std::setw(8) << "k值" << std::setw(15) << "升序(ms)" << std::setw(15) << "降序(ms)"
|
||||
<< std::setw(16) << "升序(G/s)" << std::setw(16) << "降序(G/s)" << std::endl;
|
||||
std::cout << std::string(74, '-') << std::endl;
|
||||
}
|
||||
|
||||
// 显示性能数据行
|
||||
static void printReduceSumData(int size, float time_ms, const PerformanceCalculator::ReduceSumMetrics& metrics) {
|
||||
std::cout << std::setw(12) << size << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< time_ms << std::setw(20) << std::setprecision(3) << metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
static void printSortPairData(int size, float asc_time, float desc_time,
|
||||
const PerformanceCalculator::SortPairMetrics& asc_metrics,
|
||||
const PerformanceCalculator::SortPairMetrics& desc_metrics) {
|
||||
std::cout << std::setw(12) << size << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< asc_time << std::setw(15) << desc_time << std::setw(16) << std::setprecision(3)
|
||||
<< asc_metrics.throughput_gps << std::setw(16) << desc_metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
static void printTopkPairData(int k, float asc_time, float desc_time,
|
||||
const PerformanceCalculator::TopkPairMetrics& asc_metrics,
|
||||
const PerformanceCalculator::TopkPairMetrics& desc_metrics) {
|
||||
std::cout << std::setw(8) << k << std::setw(15) << std::fixed << std::setprecision(3)
|
||||
<< asc_time << std::setw(15) << desc_time << std::setw(16) << std::setprecision(3)
|
||||
<< asc_metrics.throughput_gps << std::setw(16) << desc_metrics.throughput_gps << std::endl;
|
||||
}
|
||||
|
||||
// 显示性能文件保存消息
|
||||
static void printSavedMessage(const std::string& filename) {
|
||||
std::cout << "\n性能结果已保存到: " << filename << std::endl;
|
||||
}
|
||||
};
|
||||
|
|
@ -0,0 +1,234 @@
|
|||
#pragma once
|
||||
#include <vector>
|
||||
#include <random>
|
||||
#include <algorithm>
|
||||
#include <mc_runtime.h>
|
||||
#include <maca_fp16.h>
|
||||
#include <iostream>
|
||||
#include <chrono>
|
||||
#include <cmath>
|
||||
|
||||
// 引入模块化头文件
|
||||
#include "yaml_reporter.h"
|
||||
#include "performance_utils.h"
|
||||
|
||||
// ============================================================================
|
||||
// 测试配置常量
|
||||
// ============================================================================
|
||||
#ifndef RUN_FULL_TEST
|
||||
const int TEST_SIZES[] = {1000000, 134217728}; // 1M, 128M, 512M, 1G
|
||||
#else
|
||||
const int TEST_SIZES[] = {1000000, 134217728, 536870912, 1073741824}; // 1M, 128M, 512M, 1G
|
||||
#endif
|
||||
|
||||
const int NUM_TEST_SIZES = sizeof(TEST_SIZES) / sizeof(TEST_SIZES[0]);
|
||||
|
||||
// 性能测试重复次数
|
||||
constexpr int WARMUP_ITERATIONS = 5;
|
||||
constexpr int BENCHMARK_ITERATIONS = 10;
|
||||
|
||||
|
||||
// ============================================================================
|
||||
// 错误检查宏
|
||||
// ============================================================================
|
||||
#define MACA_CHECK(call) \
|
||||
do { \
|
||||
mcError_t error = call; \
|
||||
if (error != mcSuccess) { \
|
||||
std::cerr << "MACA error at " << __FILE__ << ":" << __LINE__ \
|
||||
<< " - " << mcGetErrorString(error) << std::endl; \
|
||||
exit(1); \
|
||||
} \
|
||||
} while(0)
|
||||
|
||||
// ============================================================================
|
||||
// 测试数据生成器
|
||||
// ============================================================================
|
||||
class TestDataGenerator {
|
||||
private:
|
||||
std::mt19937 rng;
|
||||
|
||||
public:
|
||||
TestDataGenerator(uint32_t seed = 42) : rng(seed) {}
|
||||
|
||||
// 生成随机float数组
|
||||
std::vector<float> generateRandomFloats(int size, float min_val = -1000.0f, float max_val = 1000.0f) {
|
||||
std::vector<float> data(size);
|
||||
std::uniform_real_distribution<float> dist(min_val, max_val);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = dist(rng);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机half数组
|
||||
std::vector<half> generateRandomHalfs(int size, float min_val = -100.0f, float max_val = 100.0f) {
|
||||
std::vector<half> data(size);
|
||||
std::uniform_real_distribution<float> dist(min_val, max_val);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = __float2half(dist(rng));
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机uint32_t数组
|
||||
std::vector<uint32_t> generateRandomUint32(int size) {
|
||||
std::vector<uint32_t> data(size);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = static_cast<uint32_t>(i); // 使用索引作为值，便于验证稳定排序
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成随机int64_t数组
|
||||
std::vector<int64_t> generateRandomInt64(int size) {
|
||||
std::vector<int64_t> data(size);
|
||||
for (int i = 0; i < size; i++) {
|
||||
data[i] = static_cast<int64_t>(i);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成包含NaN和Inf的测试数据 (half版本)
|
||||
std::vector<half> generateSpecialHalfs(int size) {
|
||||
std::vector<half> data = generateRandomHalfs(size, -10.0f, 10.0f);
|
||||
if (size > 100) {
|
||||
data[10] = __float2half(NAN);
|
||||
data[20] = __float2half(INFINITY);
|
||||
data[30] = __float2half(-INFINITY);
|
||||
}
|
||||
return data;
|
||||
}
|
||||
|
||||
// 生成包含NaN和Inf的测试数据 (float版本)
|
||||
std::vector<float> generateSpecialFloats(int size) {
|
||||
std::vector<float> data = generateRandomFloats(size, -10.0f, 10.0f);
|
||||
if (size > 100) {
|
||||
data[10] = NAN;
|
||||
data[20] = INFINITY;
|
||||
data[30] = -INFINITY;
|
||||
}
|
||||
return data;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 性能测试工具
|
||||
// ============================================================================
|
||||
class PerformanceMeter {
|
||||
private:
|
||||
mcEvent_t start, stop;
|
||||
|
||||
public:
|
||||
PerformanceMeter() {
|
||||
MACA_CHECK(mcEventCreate(&start));
|
||||
MACA_CHECK(mcEventCreate(&stop));
|
||||
}
|
||||
|
||||
~PerformanceMeter() {
|
||||
mcEventDestroy(start);
|
||||
mcEventDestroy(stop);
|
||||
}
|
||||
|
||||
void startTiming() {
|
||||
MACA_CHECK(mcEventRecord(start));
|
||||
}
|
||||
|
||||
float stopTiming() {
|
||||
MACA_CHECK(mcEventRecord(stop));
|
||||
MACA_CHECK(mcEventSynchronize(stop));
|
||||
float milliseconds = 0;
|
||||
MACA_CHECK(mcEventElapsedTime(&milliseconds, start, stop));
|
||||
return milliseconds;
|
||||
}
|
||||
};
|
||||
|
||||
// ============================================================================
|
||||
// 正确性验证工具
|
||||
// ============================================================================
|
||||
template<typename T>
|
||||
bool compareArrays(const std::vector<T>& a, const std::vector<T>& b, double tolerance = 1e-6) {
|
||||
if (a.size() != b.size()) return false;
|
||||
|
||||
for (size_t i = 0; i < a.size(); i++) {
|
||||
if constexpr (std::is_same_v<T, half>) {
|
||||
float fa = __half2float(a[i]);
|
||||
float fb = __half2float(b[i]);
|
||||
if (std::isnan(fa) && std::isnan(fb)) continue;
|
||||
if (std::isinf(fa) && std::isinf(fb) && (fa > 0) == (fb > 0)) continue;
|
||||
if (std::abs(fa - fb) > tolerance) return false;
|
||||
} else if constexpr (std::is_floating_point_v<T>) {
|
||||
if (std::isnan(a[i]) && std::isnan(b[i])) continue;
|
||||
if (std::isinf(a[i]) && std::isinf(b[i]) && (a[i] > 0) == (b[i] > 0)) continue;
|
||||
if (std::abs(a[i] - b[i]) > tolerance) return false;
|
||||
} else {
|
||||
if (a[i] != b[i]) return false;
|
||||
}
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
// CPU参考实现 - 稳定排序
|
||||
template<typename KeyType, typename ValueType>
|
||||
void cpuSortPair(std::vector<KeyType>& keys, std::vector<ValueType>& values, bool descending) {
|
||||
std::vector<std::pair<KeyType, ValueType>> pairs;
|
||||
for (size_t i = 0; i < keys.size(); i++) {
|
||||
pairs.emplace_back(keys[i], values[i]);
|
||||
}
|
||||
|
||||
if (descending) {
|
||||
std::stable_sort(pairs.begin(), pairs.end(),
|
||||
[](const auto& a, const auto& b) { return a.first > b.first; });
|
||||
} else {
|
||||
std::stable_sort(pairs.begin(), pairs.end());
|
||||
}
|
||||
|
||||
for (size_t i = 0; i < pairs.size(); i++) {
|
||||
keys[i] = pairs[i].first;
|
||||
values[i] = pairs[i].second;
|
||||
}
|
||||
}
|
||||
|
||||
// CPU参考实现 - TopK
|
||||
template<typename KeyType, typename ValueType>
|
||||
void cpuTopkPair(const std::vector<KeyType>& keys_in, const std::vector<ValueType>& values_in,
|
||||
std::vector<KeyType>& keys_out, std::vector<ValueType>& values_out,
|
||||
int k, bool descending) {
|
||||
std::vector<std::pair<KeyType, ValueType>> pairs;
|
||||
for (size_t i = 0; i < keys_in.size(); i++) {
|
||||
pairs.emplace_back(keys_in[i], values_in[i]);
|
||||
}
|
||||
|
||||
if (descending) {
|
||||
std::stable_sort(pairs.begin(), pairs.end(),
|
||||
[](const auto& a, const auto& b) { return a.first > b.first; });
|
||||
} else {
|
||||
std::stable_sort(pairs.begin(), pairs.end());
|
||||
}
|
||||
|
||||
keys_out.resize(k);
|
||||
values_out.resize(k);
|
||||
for (int i = 0; i < k; i++) {
|
||||
keys_out[i] = pairs[i].first;
|
||||
values_out[i] = pairs[i].second;
|
||||
}
|
||||
}
|
||||
|
||||
// CPU参考实现 - ReduceSum (使用double精度)
|
||||
template<typename InputT>
|
||||
double cpuReduceSum(const std::vector<InputT>& data, double init_value) {
|
||||
double sum = init_value;
|
||||
for (const auto& val : data) {
|
||||
if constexpr (std::is_same_v<InputT, half>) {
|
||||
float f_val = __half2float(val);
|
||||
if (!std::isnan(f_val)) {
|
||||
sum += static_cast<double>(f_val);
|
||||
}
|
||||
} else {
|
||||
if (!std::isnan(val)) {
|
||||
sum += static_cast<double>(val);
|
||||
}
|
||||
}
|
||||
}
|
||||
return sum;
|
||||
}
|
||||
|
|
@ -0,0 +1,154 @@
|
|||
#pragma once
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <map>
|
||||
#include <string>
|
||||
#include <chrono>
|
||||
#include <iomanip>
|
||||
#include <sstream>
|
||||
|
||||
// ============================================================================
|
||||
// YAML性能报告生成器
|
||||
// ============================================================================
|
||||
|
||||
class YAMLPerformanceReporter {
|
||||
public:
|
||||
struct PerformanceData {
|
||||
std::string algorithm;
|
||||
std::string input_type;
|
||||
std::string output_type;
|
||||
std::string key_type;
|
||||
std::string value_type;
|
||||
std::vector<std::map<std::string, std::string>> metrics;
|
||||
};
|
||||
|
||||
// 创建性能数据条目
|
||||
static std::map<std::string, std::string> createEntry() {
|
||||
return std::map<std::string, std::string>();
|
||||
}
|
||||
|
||||
// 生成ReduceSum性能YAML
|
||||
static void generateReduceSumYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "reduce_sum_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "ReduceSum算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"ReduceSum\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " input: \"float\"\n";
|
||||
yaml_file << " output: \"float\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("throughput_gps")) << "\n";
|
||||
yaml_file << " data_type: \"" << data.at("data_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
// 生成SortPair性能YAML
|
||||
static void generateSortPairYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "sort_pair_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "SortPair算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"SortPair\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " key_type: \"float\"\n";
|
||||
yaml_file << " value_type: \"uint32_t\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " ascending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("asc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("asc_throughput_gps")) << "\n";
|
||||
yaml_file << " descending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("desc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("desc_throughput_gps")) << "\n";
|
||||
yaml_file << " key_type: \"" << data.at("key_type") << "\"\n";
|
||||
yaml_file << " value_type: \"" << data.at("value_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
// 生成TopkPair性能YAML
|
||||
static void generateTopkPairYAML(const std::vector<std::map<std::string, std::string>>& perf_data,
|
||||
const std::string& filename = "topk_pair_performance.yaml") {
|
||||
std::ofstream yaml_file(filename);
|
||||
|
||||
// 写入头部信息
|
||||
writeHeader(yaml_file, "TopkPair算法性能测试结果");
|
||||
|
||||
// 算法信息
|
||||
yaml_file << "algorithm: \"TopkPair\"\n";
|
||||
yaml_file << "data_types:\n";
|
||||
yaml_file << " key_type: \"float\"\n";
|
||||
yaml_file << " value_type: \"uint32_t\"\n";
|
||||
|
||||
// 计算公式
|
||||
yaml_file << "formulas:\n";
|
||||
yaml_file << " throughput: \"elements / time(s) / 1e9 (G/s)\"\n";
|
||||
|
||||
// 性能数据
|
||||
yaml_file << "performance_data:\n";
|
||||
for (const auto& data : perf_data) {
|
||||
yaml_file << " - data_size: " << data.at("data_size") << "\n";
|
||||
yaml_file << " k_value: " << data.at("k_value") << "\n";
|
||||
yaml_file << " ascending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("asc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("asc_throughput_gps")) << "\n";
|
||||
yaml_file << " descending:\n";
|
||||
yaml_file << " time_ms: " << formatFloat(data.at("desc_time_ms")) << "\n";
|
||||
yaml_file << " throughput_gps: " << formatFloat(data.at("desc_throughput_gps")) << "\n";
|
||||
yaml_file << " key_type: \"" << data.at("key_type") << "\"\n";
|
||||
yaml_file << " value_type: \"" << data.at("value_type") << "\"\n";
|
||||
}
|
||||
|
||||
yaml_file.close();
|
||||
}
|
||||
|
||||
private:
|
||||
// 写入YAML文件头部
|
||||
static void writeHeader(std::ofstream& file, const std::string& title) {
|
||||
file << "# " << title << "\n";
|
||||
file << "# 生成时间: ";
|
||||
|
||||
auto now = std::chrono::system_clock::now();
|
||||
auto time_t = std::chrono::system_clock::to_time_t(now);
|
||||
file << std::put_time(std::localtime(&time_t), "%Y-%m-%d %H:%M:%S");
|
||||
file << "\n\n";
|
||||
}
|
||||
|
||||
// 格式化浮点数
|
||||
static std::string formatFloat(const std::string& value) {
|
||||
try {
|
||||
double d = std::stod(value);
|
||||
std::ostringstream oss;
|
||||
oss << std::fixed << std::setprecision(6) << d;
|
||||
return oss.str();
|
||||
} catch (...) {
|
||||
return value;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
|
@ -0,0 +1,73 @@
|
|||
# mcTileLang 算子任务文档（按 Level 分级）
|
||||
## 一、前置说明
|
||||
1. 文档参考：所有任务需遵循 mcTileLang 现有文档格式（可参考 `docs/deeplearning_operators` 下的 `matmul.md`、`elementwise.md`）。
|
||||
2. 提交要求：所有交付物需提交至 mcTileLang 仓库（Gitee 地址：https://gitee.com/metax-maca/mcTileLang ），PR 需标注“Level X 任务提交”，Issue 需关联对应任务模块。
|
||||
|
||||
|
||||
## 二、Level 1：基础验证类（5分，验证文档上限划10分，每个文档只有首次验证得分）
|
||||
### 核心目标
|
||||
熟悉 mc_tilelang 基础功能，验证现有文档的可复现性。
|
||||
|
||||
### 任务内容（二选一即可）
|
||||
**注：每个文档仅限首次验证得分，重复验证同一文档不计分**
|
||||
1. **算子文档验证**
|
||||
- 选择 1 个已完成的算子文档（`matmul.md`/`elementwise.md`/`deepseek_mla.md`/`gemv.md`），按文档中的“使用示例”编写代码，确保能成功运行。
|
||||
- 记录运行过程中的关键日志（如算子执行时间、输出结果），截图证明代码可复现，将截图提交至对应算子文档的 Gitee Issue（需注明“Level 1 验证”）。
|
||||
2. **教程文档验证**
|
||||
- 选择 1 个已完成的教程文档（`auto_tuning.md`/`debug_tools_for_tilelang.md`），按教程步骤操作（如 Auto-Tuning 参数配置、Debug 工具调用）。
|
||||
- 记录操作中发现的问题（如缺失的依赖包、步骤描述模糊点），补充完整操作流程后，将“验证结果+补充建议+操作截图”提交至教程文档的 Gitee Issue。
|
||||
|
||||
|
||||
## 三、Level 2：迁移优化类（10分）
|
||||
### 核心目标
|
||||
完成 CUDA 相关文档向 MACA 适配，实现补充算子说明文档。
|
||||
|
||||
### 任务内容（二选一或任选组合即可）
|
||||
1. **文档迁移：CUDA 内容适配 MACA**
|
||||
- 选择 1 个已完成的算子文档（如 `matmul.md`/`elementwise.md`），若文档中包含 CUDA 相关描述（如 CUDA 核函数、线程块配置），将其迁移为 MACA 编程对应的逻辑。
|
||||
- 迁移要求：明确 CUDA 语法与 MACA 语法的映射关系，补充 MACA 版本后提交迁移后的文档 PR 至 `docs/deeplearning_operators`。
|
||||
2. **教程迁移：MACA环境迁移**
|
||||
- 选择在CUDA环境上进行解释和说明的文档，将其在MACA环境下进行同样的操作并将需要补充和说明的内容提交PR到对应的文档（如 `Installation.md`）。
|
||||
|
||||
|
||||
## 四、Level 3：文档开发类（20分）
|
||||
### 核心目标
|
||||
编写未完成的算子文档，补充已有文档的 MACA 深度解析。
|
||||
|
||||
### 任务内容（二选一，需独立完成核心模块）
|
||||
1. **未完成文档编写**
|
||||
- 如从 5 个未编写的算子中选择 1个（`convolution.md`/`flash_attention.md`/`flash_linear_attention.md`/`matmul_dequant.md`/`tmac_gpu.md`），按与其他已有的算子文档结构编写文档：
|
||||
- 算子功能：说明算子的应用场景（如卷积算子用于图像特征提取）、核心计算逻辑；
|
||||
- 接口参数：列出输入（如输入张量形状、数据类型）、输出（如输出张量维度）、可选参数（如卷积核大小、步长）；
|
||||
- 使用示例：编写 mc_tilelang 调用代码（需基于 MACA 环境），包含数据构造、算子调用、结果验证；
|
||||
- 性能说明：标注算子在指定设备（如 GPU 型号）下的推荐参数配置。
|
||||
- 交付物：编写完成的算子文档，提交 PR 至 `docs/deeplearning_operators`。
|
||||
2. **已有文档 MACA 解析补充**
|
||||
- 选择 1 个已完成的算子文档（如 `matmul.md`），新增“MACA 编程解析”章节：
|
||||
- 解析内容：包含算子的 MACA 内存布局设计（如张量的 `layout` 参数配置）、计算流程优化（如循环展开、数据分块）；
|
||||
- 代码示例：提供完整的 MACA 优化版算子代码（需包含注释），对比优化前后的性能差异。
|
||||
- 交付物：补充后的文档 PR，需附代码运行截图证明优化效果。
|
||||
|
||||
|
||||
## 五、Level 4：复杂开发类（50分）
|
||||
### 核心目标
|
||||
开发复杂融合算子或贡献示例代码，满足大模型推理等高阶需求。
|
||||
|
||||
### 任务内容
|
||||
**example 文件夹代码贡献**
|
||||
- 为 `mcTileLang/examples` 文件夹新增其他算子的完整使用案例。
|
||||
- 案例要求：包含数据准备、算子调用、结果可视化（如打印特征图形状、注意力权重热力图）、性能分析代码，提交代码 PR 并附运行演示截图。
|
||||
|
||||
|
||||
## 六、验证得分规则
|
||||
- **Level 1 验证文档上限**：10分（每个文档首次验证得5分，后续重复验证不计分）
|
||||
- **验证范围**：算子文档验证和教程文档验证均适用此规则
|
||||
- **计分方式**：系统会自动识别文档验证记录，同一文档多次提交仅计算首次验证得分
|
||||
|
||||
## 七、交付物清单（按 Level 汇总）
|
||||
| Level | 必交交付物 | 可选交付物 |
|
||||
|-------|------------|------------|
|
||||
| 1 | 验证截图、Issue 提交记录 | 补充的操作步骤文档 |
|
||||
| 2 | 迁移后的文档 PR、验证截图 | 详细说明文档 |
|
||||
| 3 | 未完成算子文档 PR / 补充 MACA 解析的文档 PR | 解析逻辑说明文档 |
|
||||
| 4 | example 代码 PR | 结果输出和验证的图片 |
|
||||
|
|
@ -0,0 +1,103 @@
|
|||
# TileLang 国产GPU开发项目文档（项目背景+快速上手+生态贡献）
|
||||
## 一、项目背景
|
||||
### 1. 行业痛点：国产GPU生态的核心挑战
|
||||
当前智算与通用计算领域中，GPU内核开发面临双重困境：一方面，国际主流GPU开发语言学习成本高、代码量大，开发者需兼顾性能优化与生产力，门槛较高；另一方面，国产GPU虽在硬件性能上逐步追赶，但软件生态适配相对来说还有很大的进步空间，缺乏轻量化、高效的开发工具链，导致硬件算力难以充分释放，生态协同性不足。
|
||||
|
||||
### 2. 解决方案：TileLang的技术突破
|
||||
TileLang作为TileLang社区主导的GPU内核领域专用语言，以“高效开发+性能不妥协”为核心，针对性解决行业痛点：
|
||||
- **开发效率革命**：采用Python式简洁语法，实现FlashAttention算子仅需80行代码，并保持了与官方版本持平的性能。这种代码量的大幅减少不仅降低了开发门槛，也提高了维护性和可读性。
|
||||
- **分层接口适配**：提供3个层次编程接口，覆盖从初学者到专家的全阶段需求，降低国产GPU开发门槛。
|
||||
- **多硬件兼容**：已在MACA曦云C500、英伟达H100/A100、AMD MI250/MI300X等多类GPU上验证适配，支持“cuda/hip/cpu”多目标编译，兼容性广泛。
|
||||
|
||||
### 3. 生态协同：MACA与开源社区的落地支持
|
||||
为推动TileLang在国产GPU上的实际应用，MACA（国产高性能GPU代表厂商）与开源社区联合行动：
|
||||
- **硬件适配**：MACA AI编译器团队和TileLang社区合作已提前参与该项目，探讨MACA GPU与TileLang的适配（开源仓库：[mcTileLang](https://gitee.com/metax-maca/mcTileLang)），通过MXMACA软件栈实现深度协同，核心算子性能接近国际主流产品。
|
||||
- **在线环境搭建**：在模力方舟平台提供预配置的TileLang在线体验环境，开发者无需自行搭建硬件，直接基于曦云C500（64GB显存、Intel Xeon Gold 6530）来进行开发；
|
||||
- **资源支持**：提供专属算力券降低体验成本，同时开源完整适配代码与文档，助力生态共建。
|
||||
|
||||
### 4. 项目意义
|
||||
本项目通过“语言工具+国产硬件+在线平台”的组合，打破国产GPU生态“硬件强、软件弱”的僵局，为开发者提供“开箱即用”的国产GPU开发方案，推动中国算力产业从“单点突破”转向“生态共荣”，助力智算、通用计算领域的国产化替代与创新发展。
|
||||
|
||||
## 二、快速上手：TileLang 国产 GPU 开发实践
|
||||
### 1. 环境准备：获取TileLang在线开发资源
|
||||
#### 步骤1：进入模力方舟算力市场
|
||||
访问[模力方舟](https://ai.gitee.com/compute)，点击顶部导航栏“算力市场”，进入MACA GPU资源租用页面。
|
||||
|
||||
#### 步骤2：领取TileLang专属算力券
|
||||
- 参与比赛活动，领取TileLang专属算力券；
|
||||
- 算力券可直接兑换曦云C系列GPU容器资源，用于TileLang开发体验。
|
||||
|
||||
#### 步骤3：选择TileLang镜像与配置
|
||||
- **硬件配置选择**：默认选择“曦云C500”GPU，单卡配置为64GB显存、12核Intel Xeon Gold 6530 CPU
|
||||
- **镜像选择**：在“镜像”列表中勾选“基础镜像”下的“TileLang 0.1.5”
|
||||
- **计费方式**：支持按量收费、包日/包周/包月。
|
||||
|
||||
### 2. 容器启动与TileLang验证
|
||||
#### 步骤1：启动GPU容器
|
||||
完成配置与算力券兑换后，点击“启动容器”，等待容器初始化（通常耗时1-3分钟，可在“工作台”查看进度）。
|
||||
|
||||
#### 步骤2：执行快速验证命令
|
||||
容器启动后，通过终端输入以下命令，验证TileLang环境可用性：
|
||||
```bash
|
||||
# 进入TileLang示例目录
|
||||
cd /root/mcTileLang/
|
||||
# 运行快速启动示例
|
||||
python3 ./examples/quickstart.py
|
||||
```
|
||||
|
||||
#### 步骤3：确认验证结果
|
||||
若终端输出以下信息，说明环境正常：
|
||||
- 张量输出（如`tensor([[ -1.4619, -19.9844, ... ]], device='cuda:0', dtype=torch.float16)`）；
|
||||
- 匹配提示（`Kernel output matches PyTorch reference`）；
|
||||
- 延迟数据（如`Latency:0.11110399663448334 ms`）。
|
||||
|
||||
可额外执行`mx-smi`命令，查看曦云C500 GPU状态（如温度、显存占用、功率等），确认硬件资源正常调用。
|
||||
|
||||
### 3. 进阶体验：TileLang算子开发
|
||||
#### 步骤1：参考官方文档
|
||||
访问MACA开源项目文档，获取算子开发教程与API说明。
|
||||
|
||||
#### 步骤2：尝试核心算子开发
|
||||
基于示例代码修改，开发自定义GPU算子（如简化版GEMM算子）。
|
||||
|
||||
#### 步骤3：参与生态贡献
|
||||
若开发的算子具备通用性，可通过[gitee](https://gitee.com/metax-maca/mcTileLang)提交PR参与TileLang国产GPU生态共建。
|
||||
|
||||
|
||||
## 三、生态贡献指南：邀您共建TileLang国产GPU生态
|
||||
为持续完善mcTileLang仓库功能、降低开发门槛，诚邀所有开发者参与贡献，重点欢迎**Issue反馈**与**Docs/Example类型PR**，评分详见[Task](Task.md)具体指南如下：
|
||||
|
||||
### 1. 积极反馈：提交Issue助力仓库优化
|
||||
当您遇到以下场景时，欢迎通过仓库“Issues”模块提交反馈，帮助团队定位问题、明确方向：
|
||||
- **功能需求**：需要新增算子示例（如卷积、稀疏注意力）、补充特定硬件适配文档；
|
||||
- **Bug反馈**：运行示例代码报错、性能不符合预期、编译过程异常（需附报错日志与环境信息）；
|
||||
- **文档疑问**：现有文档（如安装步骤、API说明）不清晰、存在疏漏；
|
||||
- **优化建议**：对内核性能、接口易用性、编译速度的改进想法。
|
||||
|
||||
> 提交Issue时，请选择对应标签（如“feat request”“bug”“doc”），描述清晰场景与需求，便于社区快速响应。
|
||||
|
||||
### 2. 主动共建：提交Docs/Example类型PR
|
||||
mcTileLang仓库的`docs`（文档）与`examples`（示例）目录是生态核心组成部分，尤其欢迎以下类型PR，共建更易用的开发资源：
|
||||
|
||||
#### （1）Docs类型PR：完善文档体系
|
||||
- **教程迁移**：将已有的基于CUDA的教程文档迁移到MACA上
|
||||
- **新增教程**：补充“TileLang布局优化实战”“MACA GPU性能调优指南”等进阶教程；
|
||||
- **更新说明**：同步API变更（如新增的接口）、修正安装步骤中的过时信息；
|
||||
- **补充案例**：在文档中添加“常见问题排查”（如容器启动失败、编译报错解决）。
|
||||
|
||||
#### （2）Example类型PR：丰富算子示例
|
||||
- **新增算子**：提交RetNet、Mamba等新兴模型的TileLang实现，或补充现有算子的优化版本（如支持不同精度或者混合精度的算子）；
|
||||
- **硬件适配**：提供算子在GPU（如曦云C500）上的适配示例与性能分析
|
||||
|
||||
#### （3）PR提交流程
|
||||
1. Fork mcTileLang仓库到个人账号；
|
||||
2. 创建专属分支（如`doc/update-install-guide`、`example/add-fp8-gemm`）；
|
||||
3. 完成修改后，提交PR并关联相关Issue（如没有可以自己提交Issue并进行关联）；
|
||||
4. 参考仓库`CONTRIBUTING.md`确保代码风格、文档格式符合规范，等待审核合并。
|
||||
|
||||
|
||||
## 四、资源汇总
|
||||
- 仓库地址：https://gitee.com/metax-maca/mcTileLang
|
||||
- 在线体验：https://ai.gitee.com/compute
|
||||
- 文档参考：《曦云系列_通用计算GPU_快速上手指南》
|
||||
- 社区交流：加入社区参与讨论（仓库README有入口）
|
||||
|
|
@ -1,178 +0,0 @@
|
|||
# mcTileLang 源码编译教程
|
||||
|
||||
本教程将指导你完成 mcTileLang 的源码编译安装。
|
||||
|
||||
> **环境说明**
|
||||
> - 硬件：曦云 C500
|
||||
> - PyTorch：2.8
|
||||
> - Python：3.12
|
||||
> - Maca：3.3.0.4
|
||||
|
||||
---
|
||||
|
||||
## 0. 必选：安装 cmake
|
||||
|
||||
```bash
|
||||
apt-get update && apt-get install -y cmake
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 1. 安装依赖
|
||||
|
||||
### 1.1 基础依赖包
|
||||
|
||||
```bash
|
||||
pip install z3-solver>=4.13.0 cython psutil cloudpickle tqdm
|
||||
```
|
||||
|
||||
### 1.2 从 Maca PyPI 安装 PyTorch 相关包
|
||||
|
||||
```bash
|
||||
pip install apex dropout_layer_norm flash_attn fused_dense_lib \
|
||||
rotary_emb torch torchaudio torchvision triton \
|
||||
xentropy_cuda_lib xformers \
|
||||
-i https://repos.metax-tech.com/r/maca-pypi/simple \
|
||||
--trusted-host repos.metax-tech.com
|
||||
```
|
||||
|
||||
### 1.3 安装 tvm-ffi 依赖
|
||||
|
||||
```bash
|
||||
pip install torch-c-dlpack-ext
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 2. 从源码构建
|
||||
|
||||
### 2.1 克隆代码
|
||||
|
||||
```bash
|
||||
git clone https://gitee.com/metax-maca/mcTileLang/
|
||||
cd mcTileLang
|
||||
```
|
||||
|
||||
### 2.2 初始化子模块
|
||||
|
||||
```bash
|
||||
git submodule update --init --recursive
|
||||
```
|
||||
|
||||
> 因为众所周知的原因访问github会比较缓慢，大家如失败可以多尝试几次，实在不行可参考如下方式配置一下 `.gitmodules` 并从gitee的mirror仓库下载所需的submodule。
|
||||
|
||||
```
|
||||
[submodule "3rdparty/cutlass"]
|
||||
path = 3rdparty/cutlass
|
||||
url = https://gitee.com/topshare_test/cutlass.git
|
||||
[submodule "3rdparty/tvm"]
|
||||
path = 3rdparty/tvm
|
||||
url = https://gitee.com/topshare_test/tvm.git
|
||||
[submodule "3rdparty/composable_kernel"]
|
||||
path = 3rdparty/composable_kernel
|
||||
url = https://gitee.com/topshare_test/composable_kernel.git
|
||||
```
|
||||
|
||||
### 2.3 配置 git（首次使用需要）
|
||||
|
||||
```bash
|
||||
git config --global user.email "you@example.com"
|
||||
git config --global user.name "Your Name"
|
||||
```
|
||||
|
||||
### 2.4 必备构建工具
|
||||
|
||||
```bash
|
||||
apt-get install -y cmake patchelf
|
||||
```
|
||||
|
||||
### 2.5 CMake 配置
|
||||
|
||||
```bash
|
||||
USE_MACA=ON cmake -B build
|
||||
```
|
||||
|
||||
### 2.6 编译
|
||||
|
||||
```bash
|
||||
make -C build -j 32
|
||||
```
|
||||
|
||||
> 注：`-j 32` 表示使用 32 个线程并行编译，可根据 CPU 核心数调整
|
||||
|
||||
### 2.7 安装 tvm-ffi
|
||||
|
||||
```bash
|
||||
cd 3rdparty/tvm/3rdparty/tvm-ffi && pip install . && cd -
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 3. 运行验证
|
||||
|
||||
### 3.1 添加 tilelang 到 PYTHONPATH
|
||||
|
||||
```bash
|
||||
export PYTHONPATH=/path/to/mcTileLang
|
||||
```
|
||||
|
||||
将 `/path/to/mcTileLang` 替换为实际的 mcTileLang 路径，例如：
|
||||
|
||||
```bash
|
||||
export PYTHONPATH=/root/mcTileLang:.
|
||||
```
|
||||
|
||||
> **注意**：有时候可能需要在路径后加上 `:.` 来确保当前目录也在 Python 的搜索路径中。
|
||||
|
||||
### 3.2 验证导入
|
||||
|
||||
```bash
|
||||
python -c "import tilelang; print(tilelang.__version__)"
|
||||
```
|
||||
|
||||
### 3.3 运行简单示例
|
||||
|
||||
```bash
|
||||
python examples/quickstart.py
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 4. 任务提交
|
||||
|
||||
完成安装后，请提交以下材料：
|
||||
|
||||
1. **安装成功截图**：运行 `python -c "import tilelang; print(tilelang.__version__)"` 的输出截图
|
||||
2. **示例运行截图**：运行 `python examples/quickstart.py` 的输出截图
|
||||
3. **遇到的问题与解决方案**（如有）：记录安装过程中遇到的问题及解决方法
|
||||
|
||||
将上述材料打包提交到本仓库的 `spring_tilelang/submissions/YOUR_ID/` 目录下。
|
||||
|
||||
提交目录结构示例：
|
||||
```
|
||||
spring_tilelang/
|
||||
└── submissions/
|
||||
└── YOUR_ID/
|
||||
├── screenshots/ # 安装成功截图
|
||||
├── notes.md # 学习文档（可选）
|
||||
└── README.md # 说明文档
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 常见问题
|
||||
|
||||
| 问题 | 解决方案 |
|
||||
|------|---------|
|
||||
| `cmake: command not found` | `apt-get install -y cmake` |
|
||||
| `patchelf not found` | `apt-get install -y patchelf` |
|
||||
| git 身份验证错误 | 配置 `git config --global user.email` 和 `user.name` |
|
||||
| pip 权限警告 | 可忽略，或使用虚拟环境 |
|
||||
|
||||
---
|
||||
|
||||
## 参考资源
|
||||
|
||||
- [mcTileLang 官方仓库](https://gitee.com/metax-maca/mcTileLang)
|
||||
- [mcTileLang 文档](https://gitee.com/metax-maca/mcTileLang/tree/main/docs)
|
||||
- [TileLang 官方仓库](https://github.com/tile-ai/tilelang)
|
||||
|
|
@ -1,166 +0,0 @@
|
|||
# TileLang Puzzle 练习指南
|
||||
|
||||
本指南将帮助你通过 TileLang Puzzle 练习，掌握 TileLang 编程基础。
|
||||
|
||||
---
|
||||
|
||||
## 什么是 TileLang Puzzle？
|
||||
|
||||
TileLang Puzzles 是一系列编程练习，旨在帮助你学习 [TileLang](https://github.com/tile-ai/tilelang)——一种用于开发高性能深度学习算子的领域特定语言（DSL）。
|
||||
|
||||
练习从简单的示例开始，逐步过渡到现代算子如 GEMM 和 FlashAttention，帮助你全面理解 TileLang 的设计原理。
|
||||
|
||||
---
|
||||
|
||||
## 环境准备
|
||||
|
||||
### 前置要求
|
||||
|
||||
确保你已经完成了 [mcTileLang 源码安装](./mctilelang-install.md)。
|
||||
|
||||
### 验证环境
|
||||
|
||||
运行以下命令检查 TileLang 是否安装正确：
|
||||
|
||||
```bash
|
||||
python -c "import tilelang; print(tilelang.__version__)"
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 获取 Puzzle 代码
|
||||
|
||||
### 克隆仓库
|
||||
|
||||
```bash
|
||||
git clone https://github.com/tile-ai/tilelang-puzzles.git
|
||||
cd tilelang-puzzles
|
||||
```
|
||||
|
||||
### 目录结构
|
||||
|
||||
```
|
||||
tilelang-puzzles/
|
||||
├── puzzles/ # Puzzle 题目（需要你来完成）
|
||||
├── ans/ # 参考答案
|
||||
├── scripts/ # 辅助脚本
|
||||
└── README.md
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## Puzzle 列表
|
||||
|
||||
共有 10 个 Puzzle，难度递增：
|
||||
|
||||
| 编号 | 名称 | 描述 | 难度 |
|
||||
|------|------|------|------|
|
||||
| 01 | copy | 张量拷贝 | ⭐ |
|
||||
| 02 | vector-add | 向量加法 | ⭐ |
|
||||
| 03 | outer-vec-add | 外积向量加法 | ⭐ |
|
||||
| 04 | backward-op | 反向传播算子 | ⭐ |
|
||||
| 05 | reduce-sum | 归约求和 | ⭐ |
|
||||
| 06 | softmax | Softmax 计算 | ⭐⭐ |
|
||||
| 07 | scalar-flash-attn | 标量 Flash Attention | ⭐⭐ |
|
||||
| 08 | matrix | 矩阵运算 | ⭐⭐ |
|
||||
| 09 | conv | 卷积运算 | ⭐⭐ |
|
||||
| 10 | dequant-mm | 反量化矩阵乘法 | ⭐⭐⭐ |
|
||||
|
||||
---
|
||||
|
||||
## 如何练习
|
||||
|
||||
### 1. 运行 Puzzle
|
||||
|
||||
每个 Puzzle 都是一个独立的可执行脚本：
|
||||
|
||||
```bash
|
||||
# 运行 Puzzle 01
|
||||
python3 puzzles/01-copy.py
|
||||
```
|
||||
|
||||
### 2. 查看参考答案
|
||||
|
||||
如果遇到困难，可以参考 `ans/` 目录下的实现：
|
||||
|
||||
```bash
|
||||
# 查看 Puzzle 01 的参考答案
|
||||
python3 ans/01-copy.py
|
||||
```
|
||||
|
||||
### 3. 理解代码
|
||||
|
||||
每个 Puzzle 文件里面包含：
|
||||
- **问题描述**：注释中说明需要实现的功能
|
||||
- **待填充代码**：标记为 `# TODO` 的部分
|
||||
- **测试代码**：验证你的实现是否正确
|
||||
|
||||
---
|
||||
|
||||
## 提交方式
|
||||
|
||||
将完成的 Puzzle 提交到本仓库的 `spring_tilelang/submissions/YOUR_ID/` 目录下。
|
||||
|
||||
提交目录结构示例：
|
||||
```
|
||||
spring_tilelang/
|
||||
└── submissions/
|
||||
└── YOUR_ID/
|
||||
├── puzzles/ # 完成的 Puzzle 代码
|
||||
├── screenshots/ # 运行截图
|
||||
└── notes.md # 学习笔记
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 学习建议
|
||||
|
||||
### 循序渐进
|
||||
|
||||
1. **从简单开始**：先完成 01-03 号 Puzzle，熟悉 TileLang 语法
|
||||
2. **理解原理**：不要直接看答案，先自己思考实现方案
|
||||
3. **对比学习**：完成后再对比参考答案，学习更优写法
|
||||
4. **记录笔记**：记录每个 Puzzle 的核心概念和技巧
|
||||
|
||||
### 核心概念
|
||||
|
||||
在练习过程中，重点关注以下概念：
|
||||
|
||||
- **Tile**：TileLang 的核心抽象，表示数据块
|
||||
- **Kernel**：GPU 上执行的函数
|
||||
- **Schedule**：调度策略，影响性能
|
||||
- **Memory Layout**：内存布局，影响数据访问效率
|
||||
|
||||
### 调试技巧
|
||||
|
||||
- 使用 `print` 输出中间结果
|
||||
- 从小规模数据开始测试
|
||||
- 对比参考答案的输出来定位问题
|
||||
|
||||
---
|
||||
|
||||
## 参考资源
|
||||
|
||||
- [TileLang 官方仓库](https://github.com/tile-ai/tilelang)
|
||||
- [TileLang Puzzles 仓库](https://github.com/tile-ai/tilelang-puzzles)
|
||||
- [Triton Puzzles](https://github.com/srush/Triton-Puzzles) - 类似的 Triton 练习
|
||||
- [LeetGPU](https://leetgpu.com/) - GPU 编程练习平台
|
||||
|
||||
---
|
||||
|
||||
## 常见问题
|
||||
|
||||
### Q: 运行 Puzzle 时报错怎么办？
|
||||
|
||||
A: 首先检查 TileLang 是否正确安装，然后检查 GPU 是否可用。可以运行 `python3 scripts/check_tilelang_env.py` 检查环境。
|
||||
|
||||
### Q: 如何优化我的实现？
|
||||
|
||||
A: 可以尝试：
|
||||
- 调整 Tile 大小
|
||||
- 优化内存访问模式
|
||||
- 使用更高效的算法
|
||||
|
||||
---
|
||||
|
||||
祝你练习愉快，早日掌握 TileLang！
|
||||
|
|
@ -1,219 +0,0 @@
|
|||
# 学习文档撰写说明
|
||||
|
||||
本文档指导你如何撰写 TileLang 学习文档。
|
||||
|
||||
---
|
||||
|
||||
## 写作形式
|
||||
|
||||
你可以根据自己的学习风格选择以下任一形式：
|
||||
|
||||
### 形式一：每个 Puzzle 一篇笔记
|
||||
|
||||
为每个完成的 Puzzle 撰写单独的笔记，记录：
|
||||
- Puzzle 的核心知识点
|
||||
- 解题思路和关键代码
|
||||
- 遇到的问题及解决方法
|
||||
|
||||
**适合人群：** 喜欢详细记录、循序渐进的学习者
|
||||
|
||||
### 形式二：整体学习心得
|
||||
|
||||
完成所有练习后，撰写一篇综合性的学习总结：
|
||||
- TileLang 整体学习感悟
|
||||
- 核心概念的理解
|
||||
- 从入门到进阶的心路历程
|
||||
|
||||
**适合人群：** 喜欢宏观思考、总结提炼的学习者
|
||||
|
||||
### 形式三：混合形式
|
||||
|
||||
结合以上两种形式：
|
||||
- 简单 Puzzle 合并记录
|
||||
- 复杂 Puzzle 单独深入分析
|
||||
- 最后附整体学习感悟
|
||||
|
||||
**适合人群：** 灵活多变、按需记录的学习者
|
||||
|
||||
---
|
||||
|
||||
## 内容建议
|
||||
|
||||
### 基础内容
|
||||
|
||||
无论你选择哪种形式，建议包含以下内容：
|
||||
|
||||
1. **学习背景**
|
||||
- 为什么学习 TileLang
|
||||
- 学习目标
|
||||
|
||||
2. **核心概念理解**
|
||||
- Tile 的概念和作用
|
||||
- Kernel 的编写方式
|
||||
- Schedule 的调度策略
|
||||
|
||||
3. **实践记录**
|
||||
- 完成的 Puzzle 列表
|
||||
- 关键代码片段
|
||||
- 运行结果截图
|
||||
|
||||
4. **问题与解决**
|
||||
- 遇到的困难
|
||||
- 解决过程
|
||||
- 经验总结
|
||||
|
||||
### 进阶内容（可选）
|
||||
|
||||
- **性能分析**：对比不同实现的性能差异
|
||||
- **源码阅读**：阅读 TileLang 源码的心得
|
||||
- **拓展思考**：TileLang 与其他框架的对比
|
||||
|
||||
---
|
||||
|
||||
## 文档格式
|
||||
|
||||
### 文件命名
|
||||
|
||||
建议使用英文小写，单词间用 `-` 连接：
|
||||
|
||||
```
|
||||
puzzle-01-notes.md # 单个 Puzzle 笔记
|
||||
tilelang-learning-journey.md # 整体学习心得
|
||||
my-tilelang-notes.md # 个人笔记
|
||||
```
|
||||
|
||||
### 文档结构示例
|
||||
|
||||
#### 单个 Puzzle 笔记示例
|
||||
|
||||
```markdown
|
||||
# Puzzle 01 - Copy 学习笔记
|
||||
|
||||
## 学习目标
|
||||
理解 TileLang 中最基本的张量拷贝操作。
|
||||
|
||||
## 核心概念
|
||||
|
||||
### Tile 的概念
|
||||
Tile 是 TileLang 中数据管理的基本单位...
|
||||
|
||||
### Copy 操作的实现
|
||||
```python
|
||||
# 关键代码
|
||||
```
|
||||
|
||||
## 解题过程
|
||||
|
||||
### 第一次尝试
|
||||
...
|
||||
|
||||
### 遇到的问题
|
||||
...
|
||||
|
||||
### 最终方案
|
||||
...
|
||||
|
||||
## 总结
|
||||
通过本 Puzzle，我理解了...
|
||||
```
|
||||
|
||||
#### 整体学习心得示例
|
||||
|
||||
```markdown
|
||||
# TileLang 春节学习心得
|
||||
|
||||
## 学习背景
|
||||
春节期间参加了 TileLang 算力入门挑战...
|
||||
|
||||
## 学习历程
|
||||
|
||||
### 第一阶段：环境搭建
|
||||
...
|
||||
|
||||
### 第二阶段：基础 Puzzle
|
||||
...
|
||||
|
||||
### 第三阶段：进阶挑战
|
||||
...
|
||||
|
||||
## 核心收获
|
||||
|
||||
### 对 TileLang 的理解
|
||||
TileLang 是一个... 它的核心优势在于...
|
||||
|
||||
### 对 GPU 编程的新认识
|
||||
...
|
||||
|
||||
## 踩坑记录
|
||||
|
||||
### 坑一：...
|
||||
解决方案：...
|
||||
|
||||
### 坑二：...
|
||||
解决方案：...
|
||||
|
||||
## 未来计划
|
||||
- 深入研究...
|
||||
- 尝试优化...
|
||||
|
||||
## 结语
|
||||
这个春节没有摆烂，收获满满！
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 提交要求
|
||||
|
||||
将学习文档提交到 `spring_tilelang/submissions/YOUR_ID/` 目录：
|
||||
|
||||
```
|
||||
spring_tilelang/
|
||||
└── submissions/
|
||||
└── YOUR_ID/
|
||||
├── notes/ # 学习文档目录
|
||||
│ ├── puzzle-01.md
|
||||
│ ├── puzzle-02.md
|
||||
│ └── summary.md
|
||||
└── ...
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 写作建议
|
||||
|
||||
### 真实记录
|
||||
- 记录真实的思考过程，不要只写正确答案
|
||||
- 失败的尝试同样有价值
|
||||
- 报错信息和解决方法可以帮助他人
|
||||
|
||||
### 图文并茂
|
||||
- 适当插入代码片段
|
||||
- 添加运行截图
|
||||
- 可以画图解释概念
|
||||
|
||||
### 保持个性
|
||||
- 用自己的语言表达
|
||||
- 可以加入个人感悟
|
||||
- 风格不限，真诚即可
|
||||
|
||||
### 注重实用
|
||||
- 提供可运行的代码
|
||||
- 给出具体的命令
|
||||
- 分享实用的调试技巧
|
||||
|
||||
---
|
||||
|
||||
## 参考示例
|
||||
|
||||
以下是一些优秀的技术博客参考：
|
||||
|
||||
- [Tile-lang 介绍](https://zhuanlan.zhihu.com/p/31180917197)
|
||||
- [TileLang 官方文档](https://tilelang.com/)
|
||||
|
||||
---
|
||||
|
||||
## 最后
|
||||
|
||||
写作是最好的学习方式。通过整理和输出，你会对 TileLang 有更深入的理解。
|
||||
|
||||
祝你写作愉快，学习进步！
|
||||
Binary file not shown.
|
Before Width: | Height: | Size: 400 KiB |
Loading…
Reference in New Issue