diff --git a/agentic-kernel-optimization-on-domestic-gpu/基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md b/agentic-kernel-optimization-on-domestic-gpu/基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md
new file mode 100644
index 0000000..9a54d4e
--- /dev/null
+++ b/agentic-kernel-optimization-on-domestic-gpu/基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md
@@ -0,0 +1,198 @@
+# 题目编号：XH-202628
+
+# 基于AI Agent开发范式的国产GPU大模型推理算子库优化比赛方案
+
+## 一、发榜单位
+
+- **企业名称**：沐曦集成电路（上海）股份有限公司
+- **企业类型**：民营企业
+- **企业地址**：上海市浦东新区海科路999弄8栋C8栋。
+
+## 二、题目名称
+
+基于AI Agent开发范式的国产GPU大模型推理算子库优化
+
+## 三、题目介绍
+
+### 1. 题目背景
+
+大模型进入规模化应用阶段后，推理性能已经成为影响模型服务成本、响应速度和产业落地能力的关键因素。相比训练阶段，推理阶段具有调用频次高、并发压力大、上下文长度持续增长、部署环境复杂等特点，对底层GPU软件栈和推理算子库提出了更高要求。
+
+在大模型推理链路中，Attention、FlashAttention、KV Cache、GEMM、算子融合、量化计算、长上下文处理等底层能力，直接决定了模型推理的吞吐、延迟、显存占用和服务稳定性。以FlashInfer、FlashAttention等为代表的推理算子库，已经成为提升大模型推理效率的重要基础软件。
+
+与此同时，AI Coding与Agent技术正在改变软件开发方式。过去依赖专家经验和人工反复调参的底层优化工作，已经出现新的开发范式：通过Agent完成代码理解、性能分析、候选实现生成、自动Benchmark、错误修复、参数搜索和迭代优化，这使得Agent能力不仅是Coding Agent，而能深入GPU基础软件优化过程。
+
+### 2. 目标介绍
+
+面向国产GPU的大模型推理，参赛者需要使用/构建AI Agent，通过自动代码理解、算子迁移、性能分析、Kernel 优化、Benchmark验证等方式，优化FlashInfer、MCTLASS等推理算子库在国产GPU/MXMACA软件栈上的性能表现。
+
+以Agent/Skill驱动FlashInfer、MCTLASS等大模型推理算子库在国产GPU上的自动迁移、自动优化、自动验证与性能迭代，用Agent推动算子优化效率，用Skill沉淀优化经验，用国产GPU建立推理生态壁垒。
+
+### 3. 选题意义
+
+本赛题围绕国产GPU上的大模型推理算子库优化，实现"Agent+Skill+Benchmark+国产GPU"的新型技术路线，推动国产GPU推理生态走向"推理高性能、优化可复现、经验可沉淀、生态可持续"。
+
+## 四、参赛对象
+
+参赛对象为2026年6月1日以前正式注册的国内全日制非成人教育的普通高等学校在校专科生、本科生、硕士和博士研究生（不含在职研究生），以及全日制职业教育本科、高职高专在校学生，可通过学生赛道申报作品参赛。
+
+参赛对象可以团队或个人形式参赛，每个团队不超过10人，每件作品可由不超过3名指导教师进行指导。可以跨专业、跨学校、跨地域组队，但同一团队所有成员均应符合本赛道相关年龄、身份要求。每件作品只可由1所高等院校、科研院所等作为参赛主体提交申报。
+
+## 五、答题要求
+
+本赛题面向国产GPU大模型推理场景，鼓励参赛团队构建或使用AI Agent/Skill工作流，对FlashInfer、FlashAttention、MCTLASS（本次聚焦Fused MOE算子）等大模型推理算子库开展迁移适配、性能分析、Kernel优化、自动调优、算子融合和Benchmark验证。
+
+本赛题要求参赛团队将Agent真正用于底层算子优化过程，使其能够参与并辅助完成代码理解、性能分析、优化生成、测试验证和多轮迭代。参赛作品应体现"Agent驱动算子优化"的技术特征，而不是仅停留在概念说明、文档生成或简单代码补全层面。
+
+参赛团队需围绕国产GPU平台上的大模型推理算子库优化，完成一类或多类任务。鼓励参赛团队在保证正确性、稳定性和可复现性的前提下，提升相关算子或端到端推理链路在国产GPU上的执行性能。
+
+本赛题鼓励参赛团队使用各类AI Agent、Skill、AI Coding、智能体开发平台和自动化调优工具参与推理算子库优化过程，参赛团队可使用Claude Code、OpenCode、OpenAI Codex、OpenClaw、CodeBuddy、Cursor等开发工具，也可使用自研Agent/Skill工作流。
+
+通过本赛题，期望形成一批面向国产GPU的大模型推理算子优化成果、Agent优化工作流、Skill模块、Benchmark方法和工程实践案例，为国产GPU大模型推理生态建设提供支撑。
+
+### 任务：算子性能调优Agent/Skill
+
+| 内容 | 说明 |
+|------|------|
+| 迁移FlashInfer的关键算子到MACA平台，并优化性能 | 版本：FlashInfer-ai/FlashInfer 0.2.8
API：
BatchPrefillWithRaggedKVCacheWrapper
- headdim64/128/256, qk192+vo128 for mla
BatchPrefillWithPagedKVCacheWrapper
- headdim64/128/256
BatchMLAPagedAttentionWrapper
- headdim qk576+vo512
BatchDecodeWithPagedKVCacheWrapper
- headdim64/128/256
语言：MACA/C++
数据类型：BF16
page size：[1, 16]
seqlen：1K~180K，选自真实大模型，且seqlen取值随机以增加泛化性 |
+| 迁移FlashAttention的关键算子到MACA平台，并优化性能 | 版本：Dao-AILab/FlashAttention 2.6.3
API：flash_attn_with_kvcache
语言：MACA/C++
数据类型：BF16
headdim：[32, 64, 96, 128, 160, 192, 224, 256, 512]，其中128、256、512高优
page size：16
seqlen：1K~180K，选自真实大模型，且seqlen取值随机以增加泛化性 |
+| 迁移Fused MOE算子到MACA平台，并优化性能 | API：Fused MOE
语言：Tilelang/Triton/MACA C；
精度：INT8 W8A8；
shape: n_tokens, n_experts, topK, N, K均选自真实大模型，在典型切分，典型seqlen, batchsize的真实取值 |
+
+提交可复现的算子优化的源码、测试及测试框架、性能测试脚本、性能报告、Agent/Skill，以及对应的PPT和文档。
+
+## 六、作品评选标准
+
+本赛题采用"自动评测打榜 + 专家综合评审 + 终审擂台赛答辩"相结合的评价机制。自动评测结果作为作品性能表现的重要依据，专家评审重点考察作品的技术路线、Agent 参与深度、工程质量、创新性和可复现性。
+
+### （一）打榜机制
+
+本赛题设置阶段性排行榜和最终排行榜。参赛团队可在规定时间内多次提交作品，系统根据统一评测环境、统一数据集、统一 Benchmark 脚本和统一评分规则进行自动评测，并生成榜单排名。
+
+#### 1. 榜单设置
+
+主办方设置以下榜单：
+
+| 榜单名称 | 评价重点 | 适用作品 |
+|----------|----------|----------|
+| Agent自动优化算子性能榜 | 单个算子或算子库在指定输入规模下的性能提升，也包含Agent/Skill在优化流程中的参与深度和自动化能力（包含了Agent工具链、AutoTune、Benchmark自动化作品） | FlashInfer、FlashAttention、MCTLASS等算子优化作品 |
+
+Agent自动优化算子性能榜依据客观性能指标排名（后续将根据提交的Agent流程核对复现）。
+
+#### 2. 自动评测流程
+
+每次打榜提交后，系统按照以下流程进行评测：
+
+作品提交 → 环境检查 → 编译构建 → 正确性测试 → Benchmark性能评测 → 稳定性与回归检查 → 生成评测报告 → 更新排行榜。
+
+只有通过正确性测试和基础稳定性测试的作品，才进入性能排名。未通过正确性测试的作品不计入性能榜。
+
+#### 3. 评测指标
+
+赛题将提供对应算子的自动化测试程序和目标测试集。目标测试集中的测试用例来源于真实公开大模型或主流大模型推理负载，例如DeepSeek-V4、GLM-5.1、Kimi-K2.5、Step-3.5-Flash、MiniMax M2.5/M2.7、MiMo-V2等模型，在典型序列长度（sequence length，简称seqlen）、批大小（batch size）、head数、head dimension、数据类型等条件下的实际输入形态。
+
+测试程序将自动评测参赛作品的正确性与综合性能。其中，正确性为必须达成的基础目标，包括功能正确性、数值精度达标、边界条件通过和运行流程完整。未通过正确性测试的作品不进入性能排名。
+
+在正确性通过的前提下，测试程序将自动统计作品在目标测试集上的综合性能表现，包括延迟、吞吐、显存占用、稳定性以及不同输入规模下的覆盖情况，并按照预先公布的权重计算加权平均性能得分。该得分作为算子性能榜和相关任务排名的主要依据。
+
+#### 4. 复现与反作弊要求
+
+为保证榜单结果真实有效，参赛作品需满足以下要求：
+
+1. 作品必须能够在主办方指定环境中复现；
+2. 不得通过识别测试样例、硬编码输出、跳过计算等方式刷榜；
+3. 不得牺牲计算正确性换取性能提升；
+4. 不得利用未授权接口、系统漏洞或评测环境漏洞获取不当成绩；
+5. 主办方有权对异常成绩进行复测、要求补充说明或取消排名。
+
+如参赛作品在复现过程中无法稳定达到提交成绩，主办方可根据复测结果调整其榜单成绩。
+
+### （二）作品评选标准
+
+提交作品最终评价采用100分制，权重计划如下：
+
+| 类别 | 评审维度 | 权重 | 说明 |
+|------|----------|------|------|
+| 客观评测 | 性能提升效果 | 60% | 在保证准确性和稳定性及精度等的前提下，相比基线版本，在延迟、吞吐、Token/s、显存占用等方面取得的提升 |
+| 客观评测 | Agent/Skill 的可复现性 | 20% | Agent真实参与源码理解、代码生成、性能分析、自动调优、Benchmark和多轮迭代，按照可复现程度打分：
功能能复现：5分
性能复现达提交标称的60%以上：10分
性能复现达提交标称的80%以上：15分
性能复现达提交标称的90%以上：20分 |
+| 客观评测+主观评测 | 文档说明与演示报告 | 20% | 技术报告、README、运行说明、演示视频、答辩材料是否清晰完整，按照材料完备和质量程度打分。 |
+
+## 七、作品提交时间
+
+2026年5月至9月上旬，各参赛团队选择榜单中的题目开展研发攻关，各高校、科研机构等组织协调机构应组织学生和青年科技工作者参赛，安排专业人员给予指导，为参赛团队提供支持保障。
+
+2026年9月5日前，各参赛团队要向发榜单位完成作品提交，具体要求详见作品提交方式。
+
+2026年9月20日前，由发榜单位完成初审，确定入围终审擂台赛的晋级作品和团队。
+
+2026年10月，安排专门团队提供帮助和指导，各晋级团队完善作品。
+
+2026年11月，组织终审擂台赛，角逐"擂主"。
+
+## 八、参赛报名及作品提交方式
+
+### （一）报名方式
+
+1. 参赛选手登录"挑战杯"官网www.tiaozhanbei.net，在"揭榜挂帅"擂台赛报名入口注册账号，登录大赛申报系统在线填写报名信息。报名信息提交后，下载打印系统生成的报名表。
+2. 申报人在报名表对应位置加盖所在学校或所在单位公章。
+3. 将盖章版报名表扫描件上传至报名系统，等待系统审核。请参赛选手注意查看审核状态，如审核不通过，需重新提交。
+4. 系统开放报名时间为2026年5月30日—6月30日，逾期后系统将自动关闭报名功能。
+
+### （二）作品提交方式
+
+作品的提交除提到的客观评测外，在初赛/总决赛截止时间点前，参赛团队应将所有要求的材料，包括可复现的算子优化的源码、测试及测试框架、性能测试脚本、性能报告、Agent/Skill，以及对应的PPT和文档等统一打包压缩提交至邮箱：opensource@metax-tech.com。压缩包命名方式为：申报人所在单位-申报人姓名-作品名称-联系电话（例如：XX 大学-张XX-XX方案-手机号)。
+
+除参赛报名表外，各参赛组提交的文档、源代码和相关文件不得携带任何参赛学校、老师和学生的个人信息。同时，各参赛团队在提交作品时，同步报送1份经报名系统审核通过的参赛报名表，报名表所有信息须与系统内填报内容完全一致。
+
+## 九、赛事保障
+
+1. **算力资源支持**：参赛选手在完成报名后，提供对应线上的曦云C500等在线算力资源券。
+2. **技术培训**：赛前赛中至少组织 2 场线上培训，提供培训回放及答疑文档；另外届时会根据实际情况，决定是否组织线下的专场培训。
+3. **专家指导**：建立线上答疑社群，由沐曦股份技术团队指导，定期回复技术问题。
+4. **交流平台**：在沐曦股份开发者社区开设赛事专属社区板块，支持参赛团队分享经验、交流问题，促进技术共创。
+5. **技术文档和课程**
+
+## 十、设奖情况及奖励措施
+
+### 1. 设奖情况
+
+按参赛作品数量比例设奖，原则上评出 "擂主"1个、特等奖5个，一等奖5个、二等奖6个、三等奖8个，获奖比例不超过参赛作品总数的30%，从特等奖中角逐出擂主团队。
+
+### 2. 奖励措施
+
+- **奖金奖励**：擂主奖励10万元/个（叠加特等奖后奖金），特等奖奖励2万/个，一等奖奖励1万元/个，二等奖奖励0.5万元/个，三等奖奖励0.2万元/个。以上奖金均为税后金额。
+- **高端GPU奖励**：擂主团队2张GPU加速卡（不叠加特等奖GPU奖励），特等奖团队1张GPU加速卡。
+- **实习奖励**：优秀获奖者有机会参与之江&沐曦股份 "南湖之新"联合培养计划；所有获奖团队可获得赛事荣誉证书。
+- **曝光支持**：获奖作品可在沐曦股份开发者社区展示，作品可提供成果孵化与应用推广支持。
+
+备注：从特等奖中角逐出擂主团队，奖金10万元已叠加特等奖奖金，擂主不叠加特等奖GPU奖励。
+
+### 3. 奖金发放方式
+
+比赛结束后，单位比赛专班工作人员与获奖团队取得联系，填写奖金申请表，赛事终审结果公示无异议后（公示期约1个月），在30个工作日内通过银行转账一次性发放至团队负责人指定账户，上述所列奖金均为税后奖金。
+
+## 十一、比赛专班联系方式
+
+### 1. 专家指导团队
+
+- 顾问专家：董老师，联系电话：13482748718
+- 顾问专家：武老师，联系电话：18951640525
+- 顾问专家：韩老师，联系电话：18017474835
+
+负责比赛期间技术指导保障。
+
+### 2. 赛事服务团队
+
+- 联络专员：杨老师，联系电话：15201842467
+- 联络专员：章老师，联系电话：13501701786
+
+负责比赛期间组织服务及后期相关赛务协调联络。
+
+### 3. 联系时间
+
+比赛期间工作日（9:00－17:00）
+
+## 附：发榜单位简介
+
+沐曦集成电路（上海）股份有限公司（股票代码：688802.SH）成立于2020年9月，于2025年12月成功登陆科创板。总部位于上海，并在北京、南京、成都、杭州、深圳、武汉、长沙等地设立全资子公司及研发中心。作为一家专注于全栈GPU芯片及解决方案的集成电路设计企业，致力于打造世界一流的GPU芯片及计算平台，成为数字经济的算力基石。
+
+公司拥有技术完备、设计和产业化经验丰富的团队，核心成员平均拥有近20年高性能GPU产品端到端研发经验，曾主导过十多款世界主流高性能GPU产品研发及量产。目前，公司已推出全面覆盖人工智能训练和推理、通用计算、图形渲染和科学智能等场景的四大序列产品，并配套自研 MXMACA 软件栈，真正实现了"软硬协同"，满足"高能效"和"高通用性"的算力需求。2025年起，公司坚持"开放协同、自主可控"的方向，全面推进计算生态以及产业生态的建设。在计算生态层面，公司以自主研发的 MXMACA 全栈软件栈为核心，积极拥抱开源，打造自主、开放、兼容的通用计算开源生态；在产业生态层面，公司依托"1+6+X"战略布局，以数字算力底座为基，持续深耕金融、医疗健康、能源、教科研、交通、大文娱六大重点行业，同步积极探索具身智能、低空经济等新兴领域，为数字经济与新质生产力发展提供坚实算力支撑。
diff --git a/agentic-kernel-optimization-on-domestic-gpu/模力方舟Agent环境准备教程.md b/agentic-kernel-optimization-on-domestic-gpu/模力方舟Agent环境准备教程.md
new file mode 100644
index 0000000..e804d59
--- /dev/null
+++ b/agentic-kernel-optimization-on-domestic-gpu/模力方舟Agent环境准备教程.md
@@ -0,0 +1,688 @@
+# 模力方舟Agent部署准备教程
+
+预计用时：40 分钟
+
+本教程带你完成比赛开发环境准备：申请算力、进入模力方舟镜像、准备 baseline 代码、配置 Agent 工具和模力方舟 API Token，并完成一次本地 build / test / benchmark 验证。
+
+本教程截止到“本地 Agent 跑通”。OJ 适配与统一提交格式由主办方对 baseline 进行整理，有专门教学。
+
+## 一、教程定位
+
+本教程面向第一次在模力方舟算力实例中使用 Agent 开发的同学。完成后，你将拥有一个可直接进行代码阅读、编译、测试、benchmark 和小步优化的本地环境。
+
+本节重点解决五件事：
+
+1. 专属镜像与算力环境准备
+
+2. MXMACA / MACA 软件栈认识
+
+3. Agent 工具安装与配置
+
+4. 源码与项目目录准备
+
+5. 连通性与可用性测试
+
+
+基础镜像：
+
+```plaintext
+PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1
+```
+
+## 二、学习目标
+
+完成本教程后，你将能够：
+
+1. 申请并启动模力方舟算力实例。
+
+2. 进入 JupyterLab，并打开终端。
+
+3. 确认镜像内的 Python、PyTorch、MACA / MXMACA、mxcc 等环境可用。
+
+4. 将 baseline 源码放入 `/data` 工作目录。
+
+5. 配置模力方舟 API Token。
+
+6. 使用 OpenCode 作为 Agent 工具。
+
+7. 让 Agent 在本地完成 candidate 0 的 build、test、benchmark。
+
+8. 让 Agent 进行一次 candidate 1 小步优化，并记录 benchmark 对比。
+
+
+## 三、适用对象
+
+本教程适合：
+
+1. 已报名比赛、需要在统一镜像中开发 baseline 的选手。
+
+2. 希望使用 Agent 辅助阅读代码、定位入口、运行测试和做小步优化的选手。
+
+3. 对 MACA / MXMACA 环境不熟悉，但希望先把本地开发闭环跑通的选手。
+
+
+你不需要提前掌握 MACA kernel 开发。本教程的目标是先把环境和 Agent 闭环跑通。
+
+## 四、前置准备
+
+开始前，请确认你已经准备好以下内容。
+
+### 4.1 账号与算力
+
+1. 模力方舟网址：
+
+ [https://ai.gitee.com](https://ai.gitee.com/)
+
+2. 算力券申请入口：
+
+ [https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
+
+3. 推荐镜像：
+
+ ```plaintext
+ PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1
+ ```
+
+4. 推荐机器规格：
+
+
+沐曦GPU VRAM 16-32 GB即可
+
+### 4.2 购买模型资源包
+
+在兑换代金/算力券之后，点击模力方舟主页左侧导航栏“模型资源包”
+
+点击右上角“购买资源包” -> 选择其他资源包 -> 选择“沐曦-模型资源包”
+
+建议购买 25 元额度即可。
+
+### 4.3 设置API Key
+
+你需要一个模力方舟 API Key，用于让 Agent 调用模型。
+
+完成 4.2 的模型资源包购买后即可申请。
+
+申请方式：
+
+1. 点击进入模力方舟网页首页
+
+2. 从左侧导航栏进入“访问令牌”或“API Key”页面。
+
+3. 新建一个访问令牌。
+
+4. 复制并保存该 API Key
+
+
+后续教程中统一使用环境变量名：
+
+```plaintext
+MOARK_API_KEY
+```
+
+### 4.4 本地文件（可选）
+
+##### 如果还未获取baseline文件可跳过这一步先进行agent部署。
+
+请准备 baseline 压缩包或源码目录。示例目录名：
+
+```plaintext
+fusedmoe_v2.1
+```
+
+在模力方舟实例中，建议将代码放在：
+
+```plaintext
+/data/fusedmoe_v2.1
+```
+
+如果你的比赛材料中提供的是其他 baseline 名称，请以实际发放文件为准。
+
+## 五、知识预备
+
+### 5.1 什么是专属镜像
+
+专属镜像是主办方提前准备好的开发环境。它通常已经安装好 PyTorch、MACA / MXMACA 软件栈、编译器、驱动运行时和常用依赖。
+
+本教程使用的镜像是：
+
+```plaintext
+PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1
+```
+
+这意味着你不需要从零安装 PyTorch、MACA、mxcc 等底层组件。你需要做的是进入镜像、确认环境、放入源码并运行 baseline。
+
+### 5.2 什么是 MACA / MXMACA
+
+MACA / MXMACA 是本次算力环境中的核心软件栈，负责让程序在对应加速硬件上编译和运行。
+
+你会在命令中看到这些路径或工具：
+
+```plaintext
+MACA_PATH
+mxcc
+mxgpu_llvm
+LD_LIBRARY_PATH
+```
+
+常见路径可能是：
+
+```plaintext
+/opt/maca
+/opt/maca-20260318
+```
+
+实际路径以镜像内查询结果为准。
+
+### 5.3 什么是 Agent
+
+Agent 是能读代码、运行命令、修改文件并总结结果的编程助手。本教程推荐使用 OpenCode，并接入模力方舟的 MiniMax 模型。
+
+本教程中推荐模型：
+
+```plaintext
+MiniMax-M2.7
+```
+
+你也可以使用沐曦-模型资源包中支持的其他模型。注意：模型名必须以 API 返回的准确模型 id 为准，并且需要通过 `/v1/chat/completions` 实测可用。
+
+## 六、项目实践：环境准备与本地迭代
+
+### 步骤 1：创建算力实例
+
+目标：使用指定镜像创建比赛开发实例。
+
+操作：
+
+1. 打开模力方舟控制台，点击上方控制栏中的“算力市场”
+
+2. 进入算力实例创建页面，选择沐曦专区VRAM16-32GB算力容器
+
+3. 选择镜像：
+
+ ```plaintext
+ PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1
+ ```
+
+4. 点击创建实例。
+
+5. 创建完成后回到主页点击左侧导航栏进入“算力容器”界面
+
+
+### 步骤 2：进入 JupyterLab
+
+进入镜像中的开发环境。
+
+操作：
+
+1. 在实例列表中找到刚刚创建的实例。
+
+2. 点击“JupyterLab”或类似按钮。
+
+ 如图所示右侧按钮
+
+ 
+
+3. 页面打开后，点击 `Terminal` 新建终端。
+
+
+### 步骤 3：确认基础环境
+
+目标：确认 Python、PyTorch、mxcc、MACA 路径可用。
+
+在 JupyterLab Terminal 中执行：
+
+```plaintext
+pwd
+python --version
+which python
+
+```
+
+期望输出：
+
+
+
+继续检查 PyTorch：
+
+```plaintext
+python - <<'PY'
+import torch
+print("torch:", torch.__version__)
+print("cuda available:", torch.cuda.is_available())
+PY
+```
+
+期望输出：torch: 2.8.0+metax3.7.1.3
+
+检查 MACA / mxcc：
+
+```plaintext
+which mxcc || true
+find /opt -name mxcc 2>/dev/null | head
+ls /opt
+```
+
+预期结果：
+
+`which mxcc` 如果已经输出 mxcc 路径，说明环境变量已经配置好，可以继续下一步。
+
+如果 `which mxcc` 没有结果，但 `find /opt -name mxcc` 找到了类似路径：
+
+```plaintext
+/opt/maca/mxgpu_llvm/bin/mxcc
+```
+
+或：
+
+```plaintext
+/opt/maca-20260318/mxgpu_llvm/bin/mxcc
+```
+
+就设置环境变量。以 `/opt/maca` 为例：
+
+```plaintext
+export MACA_PATH=/opt/maca
+export PATH=$MACA_PATH/mxgpu_llvm/bin:$PATH
+export LD_LIBRARY_PATH=$MACA_PATH/lib:$MACA_PATH/mxgpu_llvm/lib:$LD_LIBRARY_PATH
+```
+
+预期结果：
+
+```plaintext
+which mxcc
+```
+
+能够输出 `mxcc` 路径。
+
+### 步骤 4：准备源码目录
+
+目标：将三个 baseline 放到 `/data` 下（左侧文件栏），作为后续 Agent 工作目录。
+
+进入 `/data`：
+
+```plaintext
+cd /data
+ls
+```
+
+假设上传后的文件名是：
+
+```plaintext
+fusedmoe_v2.1.zip
+```
+
+解压：
+
+```plaintext
+cd /data
+python - <<'PY'
+import zipfile
+from pathlib import Path
+
+zip_path = Path("fusedmoe_v2.1.zip")
+with zipfile.ZipFile(zip_path) as zf:
+ zf.extractall(".")
+print("extracted", zip_path)
+PY
+```
+
+说明：请把上面整段 Python 命令一次性复制到终端中运行。
+
+进入项目：
+
+```plaintext
+cd /data/fusedmoe_v2.1
+find . -maxdepth 3 -type f | sort | head -120
+```
+
+注意：
+
+解压后目录名可能会改变。如果 `cd /data/fusedmoe_v2.1` 提示目录不存在，请先执行 `ls /data` 查看实际解压出的目录名，再进入对应目录。
+
+你现在应该能看到类似入口：
+
+```plaintext
+scripts/build_fused_moe_i8_tn_pybind.sh
+scripts/run_fused_moe_i8_tn_pybind_test.sh
+scripts/run_fused_moe_i8_tn_benchmark.sh
+standalone/fused_moe_i8_tn/
+```
+
+源码目录准备完成，后续命令都在 `/data/fusedmoe_v2.1` 或实际解压出的 baseline 目录中执行。
+
+### 步骤 5：配置模力方舟 API Key
+
+目标：让后续 Agent 能调用模力方舟模型。
+
+在终端中设置：
+
+```plaintext
+echo 'export MOARK_API_KEY="你的真实 API Key"' >> ~/.bashrc
+source ~/.bashrc
+```
+
+确认环境变量存在：
+
+```plaintext
+test -n "$MOARK_API_KEY" && echo "MOARK_API_KEY is set" || echo "MOARK_API_KEY is missing"
+```
+
+期望输出：
+
+MOARK\_API\_KEY is set
+
+注意：
+
+不要把 API Key 粘贴到公开仓库、提交代码或截图中。
+
+如果新开 Terminal 后显示 `MOARK_API_KEY is missing`，说明环境变量还没有写入当前 shell。重新执行上面的 `source ~/.bashrc` 即可。
+
+如果后续更换 API Key，请先打开 `~/.bashrc` 删除旧的 `MOARK_API_KEY` 行，或执行下面命令追加新 key 后重新加载：
+
+```bash
+echo 'export MOARK_API_KEY="你的新 API Key"' >> ~/.bashrc
+source ~/.bashrc
+```
+
+如果同一个文件里出现多行 `MOARK_API_KEY`，通常最后一行会生效。为了避免混乱，建议只保留一行。
+
+### 步骤 6：测试模力方舟 API 连通性
+
+目标：确认 Token 与模型可用。
+
+测试文档中使用 `MiniMax-M2.7`：
+
+```plaintext
+curl https://api.moark.com/v1/chat/completions \
+ -H "Authorization: Bearer $MOARK_API_KEY" \
+ -H "Content-Type: application/json" \
+ -d '{
+ "model": "MiniMax-M2.7",
+ "messages": [{"role": "user", "content": "你好，请用一句话回复。"}],
+ "stream": false
+ }'
+```
+
+预期结果：
+
+返回 JSON，并且能看到模型回复内容。
+
+#### 替换模型
+
+只能替换沐曦模型资源包中支持的模型。先用以下命令查询当前 API Key 能看到的模型 id：
+
+```bash
+curl https://api.moark.com/v1/models \
+ -H "Authorization: Bearer $MOARK_API_KEY" \
+ | python -m json.tool
+```
+
+返回结果中 `"id"` 后面的字段就是模型 id。请复制完整 id，不要手动猜模型名。
+
+注意：`/v1/models` 中能看到某个模型，不代表它一定能用于 `/v1/chat/completions`。复制模型 id 后，还需要把步骤 6 第一段 curl 中的 `"model"` 字段改成该 id 并再次测试。只有 curl 实测返回正常回复的模型，才建议写入 OpenCode 配置。
+
+### 步骤 7：安装 OpenCode
+
+目标：安装并启动 Agent 工具。
+
+先检查 Node.js 和 npm：
+
+```plaintext
+node --version
+npm --version
+```
+
+如果镜像中已经安装 npm，执行：
+
+```plaintext
+npm install -g opencode-ai
+```
+
+验证：
+
+```plaintext
+opencode --version
+```
+
+### 步骤 8：配置 OpenCode 使用 Moark
+
+目标：让 OpenCode 通过模力方舟 API 调用 `MiniMax-M2.7`。
+
+创建配置目录：
+
+```plaintext
+mkdir -p ~/.config/opencode
+```
+
+写入配置（可以后续覆写添加模型）：
+
+```plaintext
+python - <<'PY'
+import json
+import os
+from pathlib import Path
+
+key = os.environ["MOARK_API_KEY"]
+
+cfg = {
+ "$schema": "https://opencode.ai/config.json",
+ "provider": {
+ "moark": {
+ "npm": "@ai-sdk/openai-compatible",
+ "name": "Moark",
+ "options": {
+ "baseURL": "https://api.moark.com/v1",
+ "apiKey": key,
+ },
+ "models": {
+ "MiniMax-M2.7": {
+ "name": "MiniMax-M2.7",
+ "limit": {
+ "context": 200000,
+ "output": 200000,
+ },
+ },
+
+ # 可选：如需添加其他模型，请先用步骤6中 curl 指令验证该模型可用，
+ # 然后取消下面示例配置的注释，并把“其他模型名”
+ # 替换为 /v1/models 返回的准确模型 id。
+ #
+ # "其他模型名": {
+ # "name": "其他模型名",
+ # "limit": {
+ # "context": 200000,
+ # "output": 200000,
+ # },
+ # },
+ },
+ }
+ },
+}
+
+config_path = Path.home() / ".config/opencode/opencode.json"
+config_path.parent.mkdir(parents=True, exist_ok=True)
+config_path.write_text(json.dumps(cfg, ensure_ascii=False, indent=2))
+
+print("wrote", config_path)
+PY
+```
+
+如需添加多个模型，可以在配置文件的 `models` 字段中添加多个下面这种模块。添加前必须先用步骤 6 的 curl 命令确认该模型可用于 `/v1/chat/completions`：
+
+```python
+"其他模型名": {
+ "name": "其他模型名",
+ "limit": {
+ "context": 200000,
+ "output": 200000,
+ },
+},
+```
+
+注意：重新运行上面的 Python 配置脚本会覆盖整个 `~/.config/opencode/opencode.json`。如果要保留多个模型，请在同一次配置中把所有模型都写进 `models` 字段。
+
+检查配置：
+
+```plaintext
+grep -n "baseURL\|MiniMax\|apiKey" ~/.config/opencode/opencode.json
+```
+
+如果配置中显示：
+
+```plaintext
+"apiKey": "${MOARK_API_KEY}"
+```
+
+说明 OpenCode 可能不会自动展开环境变量，导致 `Unauthorized`。为了先跑通，本教程推荐用上面的 Python 脚本把真实 key 写入本机配置文件。
+
+#### 如何使用其他模型：
+
+如果需要使用其他模型，请先用步骤 6 的 curl 方法把 "model" 字段改成目标模型名，确认 API 能返回结果。
+
+确认可用后，再把该模型加入 OpenCode 配置中的 models 字段。启动 OpenCode 后使用 /models 切换 Moark 下的模型。
+
+### 步骤 9：启动 OpenCode
+
+目标：进入项目目录并启动 Agent。
+
+```plaintext
+cd /data/fusedmoe_v2.1
+opencode
+```
+
+页面加载完毕后如果需要选择模型，请输入：
+
+```plaintext
+/models
+```
+
+只能选择 Moark 下面已经配置过、且 curl 实测可用的模型。其他模型可能因为资源包、接口类型或模型名不匹配而报错。
+
+进入后，先发送以下测试提示词测试运行情况：
+
+```text
+请通过真实 shell 命令检查当前目录、Python 版本、mxcc 路径、MACA_PATH 和 baseline 脚本入口。然后检查当前目录结构，并找出 fusedmoe_v2.1 的 build、test、benchmark 入口。先不要修改代码。
+```
+
+## 七、Agent 使用说明
+
+### 7.1 推荐工作方式
+
+使用 Agent 时，不要一开始就让它“直接优化到最快”。推荐节奏是：
+
+1. 先让它检查目录。
+
+2. 再让它跑通原始 baseline。
+
+3. 然后让它只做一处小修改。
+
+4. 每次修改后必须 build、test、benchmark。
+
+5. 每轮都记录结果。
+
+
+### 7.2 使用 Agent 的注意事项
+
+1. 不要相信没有命令输出支持的环境判断。
+
+2. 不要让 Agent 同时修改很多文件。
+
+3. 不要让 Agent 修改测试标准。
+
+
+## 八、常见问题
+
+### 问题 1：OpenCode 报 Unauthorized
+
+可能原因：
+
+1. `MOARK_API_KEY` 没有设置。
+
+2. OpenCode 配置没有读到真实 key。
+
+3. 配置里写的是字面量 `"${MOARK_API_KEY}"`。
+
+
+检查：
+
+```plaintext
+test -n "$MOARK_API_KEY" && echo "MOARK_API_KEY is set" || echo "MOARK_API_KEY is missing"
+grep -n "apiKey" ~/.config/opencode/opencode.json
+```
+
+解决：
+
+重新执行“步骤 8：配置 OpenCode 使用 Moark”中的 Python 配置脚本。
+
+### 问题 2：MiniMax-M2.1 不可用
+
+如果接口返回资源包不支持 `MiniMax-M2.1`，直接使用：
+
+```plaintext
+MiniMax-M2.7
+```
+
+或任意沐曦-模型资源包包含的模型
+
+本教程使用 `MiniMax-M2.7`。
+
+### 问题 3：找不到 mxcc
+
+检查：
+
+```plaintext
+which mxcc || true
+find /opt -name mxcc 2>/dev/null | head
+```
+
+如果找到了 mxcc，但 `which mxcc` 为空，说明环境变量没有设置。根据实际路径设置 `MACA_PATH`、`PATH` 和 `LD_LIBRARY_PATH`。
+
+### 问题 4：Python.h file not found
+
+先检查 Python 头文件：
+
+```plaintext
+python - <<'PY'
+import sysconfig, os
+inc = sysconfig.get_paths()["include"]
+print("python include:", inc)
+print("Python.h exists:", os.path.exists(os.path.join(inc, "Python.h")))
+print("LIBDIR:", sysconfig.get_config_var("LIBDIR"))
+print("LDLIBRARY:", sysconfig.get_config_var("LDLIBRARY"))
+PY
+```
+
+如果 `Python.h exists: True`，但 build 仍失败，请用：
+
+```plaintext
+bash -x scripts/build_fused_moe_i8_tn_pybind.sh
+```
+
+查看完整编译命令中是否带上了 Python include 路径。
+
+### 问题 5：加载到旧的 .so
+
+如果测试时报类似旧 Python 动态库错误，例如：
+
+```plaintext
+libpython3.10.so.1.0: cannot open shared object file
+```
+
+说明可能加载了旧编译产物。可以清理后重新 build：
+
+```plaintext
+cd /data/fusedmoe_v2.1
+rm -rf build standalone/fused_moe_i8_tn/build
+find . -name "*.so" -delete
+bash scripts/build_fused_moe_i8_tn_pybind.sh
+```
+
+### 问题 6：Agent 编造路径或测试结果
+
+处理方式：
+
+让 Agent 先执行真实命令：
+
+```plaintext
+请不要猜测。先运行 pwd、find . -maxdepth 3 -type f、python --version、which mxcc，并基于真实输出继续分析。
+```
\ No newline at end of file
diff --git a/agentic-kernel-optimization-on-domestic-gpu/模力方舟快速使用SOP.md b/agentic-kernel-optimization-on-domestic-gpu/模力方舟快速使用SOP.md
new file mode 100644
index 0000000..dc8fc58
--- /dev/null
+++ b/agentic-kernel-optimization-on-domestic-gpu/模力方舟快速使用SOP.md
@@ -0,0 +1,39 @@
+# 模力方舟快速使用SOP
+
+## 1.获取算力券
+
+[https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
+
+## 2.兑换算力和登陆平台
+
+[https://ai.gitee.com/](https://ai.gitee.com/)
+
+进入费用中心 - 算力券 ，点击右上角 “兑换”。
+
+
+
+## 3.租用算力
+
+进入算力容器，选择沐曦，租用算力，建议优先选16G显存/32G显存，如下图：
+
+
+
+## 4.创建实例
+
+基础镜像：maca-pytorch:3.7.1.5-torch2.8-py312-ubuntu24.04-amd64
+
+后续会有专属镜像【敬请期待】。
+
+
+
+TileLang赛道专属镜像：
+
+
+
+## 5.项目创作
+
+进入我的算力容器，刚创建的实例默认开机状态，点击工具-lab开始项目创作，不用时记得关机！
+
+
+
+小提示：也可以用本地AI coding工具ssh连接云服务器（如何实现和AI对话即可）
\ No newline at end of file
diff --git a/tilelang-llm-inference-operator-optimization/race-1-domestic-software-stack-samples.md b/tilelang-llm-inference-operator-optimization/race-1-domestic-software-stack-samples.md
new file mode 100644
index 0000000..6d6d220
--- /dev/null
+++ b/tilelang-llm-inference-operator-optimization/race-1-domestic-software-stack-samples.md
@@ -0,0 +1,19 @@
+# 揭榜挂帅第一赛题——基于国产软件栈的大模型推理前沿算子优化（TileLang）Sample 算子获取说明
+
+本文档用于统一维护揭榜挂帅第一赛题「基于国产软件栈的大模型推理前沿算子优化（TileLang）」提供的 Sample 算子获取入口
+
+
+- 源码仓库: [tile-ai/tilelang-metax](https://github.com/tile-ai/tilelang-metax)
+
+## Sample 算子获取地址
+
+| 算子名称 | 获取地址 |
+| --- | --- |
+| MLA | [race_tests/mla](https://github.com/tile-ai/tilelang-metax/tree/ee6db4376484f2f7270183c01fd0d90f794965cb/race_tests/mla) |
+| MoE | [race_tests/moe](https://github.com/tile-ai/tilelang-metax/tree/ee6db4376484f2f7270183c01fd0d90f794965cb/race_tests/moe) |
+| NSA | [race_tests/nsa](https://github.com/tile-ai/tilelang-metax/tree/ee6db4376484f2f7270183c01fd0d90f794965cb/race_tests/nsa) |
+
+## 使用说明
+
+
+如需查看赛题最新更新内容，可访问 race 分支目录：[race_tests](https://github.com/tile-ai/tilelang-metax/tree/race/race_tests)。
\ No newline at end of file
diff --git a/tilelang-llm-inference-operator-optimization/race_tests_run_guide基于tilelang算子sample跑通测试.md b/tilelang-llm-inference-operator-optimization/race_tests_run_guide基于tilelang算子sample跑通测试.md
new file mode 100644
index 0000000..b2b946d
--- /dev/null
+++ b/tilelang-llm-inference-operator-optimization/race_tests_run_guide基于tilelang算子sample跑通测试.md
@@ -0,0 +1,267 @@
+# TileLang-MetaX Race 分支 `race_tests` 三大算子测试跑通指南
+
+> 本文档记录如何在**模力方舟**（MetaX GPU）环境下，完成 `tilelang-metax` **race** 分支中 `race_tests` 目录下三个核心算子（MLA、MoE、NSA）的功能与性能测试。
+
+---
+
+## 1. 前置环境与硬件信息
+
+### 1.1 硬件与驱动
+
+```text
+MX-SMI 2.2.12
+Kernel Mode Driver Version: 3.0.11
+MACA Version: 3.5.3.20
+GPU: MetaX C500
+VRAM: 65536 MiB
+Sliced GPU: 50% Compute, 32000 MiB Vram Quota
+```
+
+> 本次测试在 **模力方舟** 平台提供的 MetaX C500 容器镜像中完成，镜像已预装 MACA 工具链与 PyTorch（`torch 2.8.0+metax3.5.3.9`）。
+
+### 1.2 软件版本
+
+| 组件 | 版本 |
+|------|------|
+| Python | 3.12.11 |
+| PyTorch | 2.8.0+metax3.5.3.9 |
+| TileLang | 0.1.9+maca.gitf1ca0fb9 |
+| CMake | 4.3.2 |
+| GCC | 11.4.0 |
+
+---
+
+## 2. 仓库准备与编译安装
+
+### 2.1 拉取代码并切换到 race 分支
+
+```bash
+cd /data
+git clone https://github.com/tile-ai/tilelang-metax.git
+cd tilelang-metax
+git checkout race
+```
+
+### 2.2 编译安装 TileLang（MACA 后端）
+
+在容器内执行以下命令，启用 MACA 后端并重新编译：
+
+```bash
+cd /app/tilelang-metax # 或 /data/tilelang-metax
+rm -rf build
+export USE_MACA=ON
+pip install -e . -v
+```
+
+编译完成后，验证安装：
+
+```bash
+python -c "import tilelang; print(tilelang.__version__)"
+# 输出：0.1.9+maca.gitf1ca0fb9
+```
+
+> 编译日志显示成功构建了 692 个目标，包括 `libtilelang.so`、`libtvm.so` 及 Cython wrapper，最终生成 wheel 并安装到当前 Python 环境。
+
+---
+
+## 3. 算子测试总览
+
+race 分支在 `race_tests/` 目录下提供了三个竞赛算子测试：
+
+| 算子 | 目录 | 说明 | 测试脚本 |
+|------|------|------|----------|
+| **MLA** | `race_tests/mla/` | Multi-Head Latent Attention（多头潜在注意力）FlashAttention 风格分块实现 | `test_tilelang_mla.py` |
+| **MoE** | `race_tests/moe/` | Mixture of Experts（混合专家模型）DeepSeek 风格 Routed Grouped GEMM | `fusedmoe_benchmark.py` |
+| **NSA** | `race_tests/nsa/` | Native Sparse Attention（原生稀疏注意力）前向推理 | `test_tilelang_nsa_fwd.py` |
+
+---
+
+## 4. MLA（Multi-Head Latent Attention）测试
+
+### 4.1 算子简介
+
+MLA 测试基于 TileLang 实现了类 FlashAttention 的 split / no_split 双路径内核，支持：
+- **Q/KV 分离**：`Q` + `Q_pe` 与 `KV` + `K_pe`
+- **分块计算**：通过 `block_N`、`block_H` 控制 shared memory 分块
+- **精度校验**：与 PyTorch 参考实现 (`ref_program`) 对比，`rtol=2e-4, atol=1e-4`
+
+### 4.2 运行方式
+
+#### 单 case 快速测试
+
+```bash
+cd race_tests/mla
+python test_tilelang_mla.py \
+ --no-json \
+ --batch 1 --heads 16 --kv_heads 1 \
+ --kv_ctx 2048 --dim 512 --pe_dim 64
+```
+
+#### 批量回归测试（JSON 用例）
+
+```bash
+cd /data/tilelang-metax
+python race_tests/mla/test_tilelang_mla.py
+```
+
+用例文件：`race_tests/mla/test_cases_mla_batch_ctx.json`，共 **31 组** 不同 batch 与 context length 的组合。
+
+### 4.3 测试结果
+
+**全部 31 组用例通过（31/31 PASS）**，结果摘要如下：
+
+```text
+=== Summary: 31/31 passed, 0/31 failed ===
+CSV saved to: race_tests/mla/test_cases_mla_batch_ctx_results.csv
+```
+
+#### 关键数据节选
+
+| case_id | batch | kv_ctx | latency_ms | tflops | status |
+|---------|-------|--------|------------|--------|--------|
+| 1 | 1 | 2048 | 1.1888 | 0.0600 | PASS |
+| 3 | 1 | 8192 | 4.5795 | 0.0623 | PASS |
+| 6 | 1 | 65536 | 36.3962 | 0.0627 | PASS |
+| 12 | 4 | 2048 | 1.3289 | 0.2146 | PASS |
+| 16 | 4 | 65536 | 41.3467 | 0.2207 | PASS |
+| 27 | 32 | 2048 | 1.3790 | 1.6546 | PASS |
+| 31 | 32 | 65536 | 43.4431 | 1.6807 | PASS |
+
+**观察**：
+- 随着 `batch` 增大，吞吐量（TFlops）线性提升，最高达到 **~1.74 TFlops**（batch=32）。
+- Latency 随 `kv_ctx` 增长而增加，但保持在合理范围内。
+
+---
+
+## 5. MoE（Mixture of Experts）测试
+
+### 5.1 算子简介
+
+MoE 测试实现了 DeepSeek 风格的 **Routed Grouped GEMM**：
+- **Gating Network**：Top-K 路由选择专家
+- **Grouped GEMM**：通过 `RoutedMoEKernel` 完成各专家的 gate/up/down 投影
+- **Scatter-Reduce**：将各专家输出按索引汇总回原始 token 位置
+
+测试包含两类：
+- **Functional**：与 `ref_fusedmoe.py` 的参考实现逐元素对比，`atol=1e-2, rtol=1e-2`
+- **Performance**：CUDA Event 计时，warmup 10 轮，迭代 100 轮取平均
+
+### 5.2 运行方式
+
+```bash
+cd race_tests/moe
+python fusedmoe_benchmark.py
+# 或
+bash run.sh
+```
+
+用例配置：`moe_test_configs.json`，包含 2 组功能测试 + 2 组性能测试。
+
+### 5.3 测试结果
+
+```text
+=== Running functional tests ===
+✅ Functional test passed for config: {'dhidden': 7168, 'dexpert': 2048, 'nroutedexperts': 8, 'nexpertspertoken': 4, 'bs': 1, 'seqlen': 8192, 'seed': 81394}
+✅ Functional test passed for config: {'dhidden': 3584, 'dexpert': 1024, 'nroutedexperts': 4, 'nexpertspertoken': 2, 'bs': 2, 'seqlen': 4096, 'seed': 81394}
+
+=== Running performance tests ===
+⏱ Performance test: 460.92968750ms for config: {'dhidden': 7168, 'dexpert': 2048, 'nroutedexperts': 8, 'nexpertspertoken': 4, 'bs': 4, 'seqlen': 8192, 'seed': 81394}
+⏱ Performance test: 66.26872559ms for config: {'dhidden': 3584, 'dexpert': 1024, 'nroutedexperts': 4, 'nexpertspertoken': 2, 'bs': 8, 'seqlen': 4096, 'seed': 81394}
+```
+
+**结论**：
+- **2/2 功能测试全部通过**，TileLang 实现与参考实现数值一致。
+- **2/2 性能测试完成**，在大 batch（bs=4, seqlen=8192, 8 专家）场景下单次前向约 **460.9 ms**；较小规模（bs=8, seqlen=4096, 4 专家）约 **66.3 ms**。
+
+---
+
+## 6. NSA（Native Sparse Attention）测试
+
+### 6.1 算子简介
+
+NSA 测试实现了**原生稀疏注意力前向推理**：
+- **Block-wise Sparse**：通过 `block_indices` 仅加载选定的 KV block，降低显存与计算量
+- **Causal Mask 支持**：`is_causal=True` 时自动屏蔽未来位置
+- **Shared Memory 预估**：测试脚本内置了 shared memory 占用估算，超出 65536 Bytes 的 case 自动 skip
+
+### 6.2 运行方式
+
+```bash
+cd race_tests/nsa
+python test_tilelang_nsa_fwd.py
+```
+
+用例文件：`test_cases_nsa_fwd.json`，共 **109 组** 参数组合，覆盖不同 batch、seq_len、head dim、block_size、selected_blocks 等。
+
+### 6.3 测试结果
+
+测试持续运行中，前 34 组 case 全部通过，典型输出如下：
+
+```text
+[1/109] B=1 SEQ_LEN=64 H=1 HQ=16 D=32 S=1 block_size=16 is_causal=True
+ GPU latency: 0.0309 ms
+[2/109] B=1 SEQ_LEN=64 H=1 HQ=16 D=64 S=1 block_size=16 is_causal=True
+ GPU latency: 0.0236 ms
+...
+[29/109] B=2 SEQ_LEN=1024 H=1 HQ=16 D=128 S=1 block_size=16 is_causal=True
+ GPU latency: 0.0891 ms
+[30/109] B=4 SEQ_LEN=64 H=1 HQ=16 D=32 S=1 block_size=16 is_causal=True
+ GPU latency: 0.0227 ms
+```
+
+结果自动追加写入：`benchmark_results_nsa_fwd.csv`
+
+#### CSV 结果节选
+
+| idx | B | SEQ_LEN | H | HQ | D | S | block_size | latency_ms | status |
+|-----|---|---------|---|----|---|---|------------|------------|--------|
+| 1 | 1 | 64 | 1 | 16 | 32 | 1 | 16 | 0.030863 | PASS |
+| 5 | 1 | 128 | 1 | 16 | 64 | 1 | 16 | 0.029292 | PASS |
+| 14 | 1 | 1024 | 1 | 16 | 64 | 1 | 16 | 0.025723 | PASS |
+| 29 | 2 | 1024 | 1 | 16 | 128 | 1 | 16 | 0.089119 | PASS |
+
+**观察**：
+- 小规模（B=1, SEQ_LEN=64~128）latency 约 **0.02~0.03 ms**
+- 中等规模（B=2, SEQ_LEN=1024, D=128）latency 约 **0.09 ms**
+- 所有已运行 case 状态均为 **PASS**，无 shared memory 超限导致的 skip
+
+---
+
+## 7. 总结
+
+| 算子 | 测试类型 | 用例数 | 通过数 | 关键结论 |
+|------|----------|--------|--------|----------|
+| **MLA** | 功能+性能 | 31 | 31/31 | FlashAttention 风格实现正确，batch=32 时达 ~1.74 TFlops |
+| **MoE** | 功能+性能 | 4 | 4/4 | Routed Grouped GEMM 数值正确，大模型配置单次前向约 460 ms |
+| **NSA** | 功能+性能 | 109 | 持续通过中 | 稀疏注意力前向推理稳定，小尺度 latency < 0.1 ms |
+
+**环境确认**：
+- ✅ MetaX C500 + MACA 3.5.3.20 驱动就绪
+- ✅ TileLang-MetaX race 分支编译安装成功
+- ✅ `race_tests` 三大算子均可在当前环境正常编译、运行并通过精度校验
+
+---
+
+## 8. 附录：常用命令速查
+
+```bash
+# 查看 GPU 状态
+mx-smi
+
+# 验证 TileLang 版本
+python -c "import tilelang; print(tilelang.__version__)"
+
+# MLA 测试
+python race_tests/mla/test_tilelang_mla.py
+
+# MoE 测试
+cd race_tests/moe && python fusedmoe_benchmark.py
+
+# NSA 测试
+cd race_tests/nsa && python test_tilelang_nsa_fwd.py
+```
+
+---
+
+*文档生成时间：2026-05-25*
+*测试平台：模力方舟 MetaX C500*
diff --git a/tilelang-llm-inference-operator-optimization/tilelang_maca_build_guide_模力方舟.md b/tilelang-llm-inference-operator-optimization/tilelang_maca_build_guide_模力方舟.md
new file mode 100644
index 0000000..e2910d2
--- /dev/null
+++ b/tilelang-llm-inference-operator-optimization/tilelang_maca_build_guide_模力方舟.md
@@ -0,0 +1,288 @@
+# TileLang on MACA 编译与导入完整指南
+
+> 记录从源码编译 tilelang-metax 并启用 MACA 后端的完整过程
+> 环境：MetaX C500 / MACA 3.5.3.20 / Python 3.12 / Conda
+
+---
+
+## 1. 项目背景
+
+**代码位置**：`/app/tilelang-metax`（沐曦适配分支，dev 分支，commit f1ca0fb9）
+
+**版本信息**：
+- 基础版本：`0.1.9`
+- Git 描述：`v0.1.9-73-gf1ca0fb9`
+- 构建标识：`+maca.gitf1ca0fb9`（启用 MACA 后）
+
+---
+
+## 2. 初始状态排查
+
+### 2.1 检查是否已安装
+
+```bash
+pip show tilelang
+# 结果：未安装
+
+python -c "import tilelang; print(tilelang.__version__)"
+# 结果：import 失败
+```
+
+### 2.2 查找源码位置
+
+```bash
+find / -name "tilelang*" -type d 2>/dev/null | head -10
+# 结果：/app/tilelang-metax
+```
+
+### 2.3 检查 GPU 环境
+
+```bash
+mx-smi
+# 结果：MetaX C500，MACA 3.5.3.20，/opt/maca 已安装
+```
+
+---
+
+## 3. 第一次编译尝试（失败）
+
+### 3.1 命令
+
+```bash
+cd /app/tilelang-metax
+pip install -e . -v
+```
+
+### 3.2 失败原因
+
+**错误信息**：
+```
+ninja: error: Makefile:5: expected '=', got ':'
+default_target: all
+ ^ near here
+```
+
+**根因**：`/app/tilelang-metax/build/` 目录下已有之前用 **Unix Makefiles** 生成的构建产物（包含 `Makefile`），但 pip 安装时使用了 **Ninja** 构建系统���Ninja 误把 `Makefile` 当成 ninja 构建文件解析，导致语法错误。
+
+---
+
+## 4. 清理后重新编译（成功，但缺少 MACA）
+
+### 4.1 命令
+
+```bash
+cd /app/tilelang-metax
+rm -rf build
+pip install -e . -v
+```
+
+### 4.2 编译结果
+
+- **构建成功**，版本号为 `0.1.9+cuda.gitf1ca0fb9`
+- **但 MACA 后端未启用**：CMake 日志显示 `-- CUDA toolkit not found; building without CUDA support by default.`
+- 未检测到 `USE_MACA` 环境变量，因此 TVM 未注册 `maca` target
+
+---
+
+## 5. 导入报错与修复
+
+### 5.1 导入失败
+
+```bash
+python -c "import tilelang"
+```
+
+**错误信息**：
+
+```
+ValueError: Target kind "maca" is not defined.
+Target creation from string failed: maca
+```
+
+**根因**：`tilelang/utils/target.py` 的 `determine_target("auto")` 检测到系统有 MACA 环境（`mxcc.find_maca_path()` 成功），于是返回 `"maca"`，但当前构建未包含 MACA 后端，`Target("maca")` 创建失败。
+
+### 5.2 修复 1：gemm_sp.py 导入保护
+
+**文件**：`tilelang/language/experimental/gemm_sp.py`
+
+**修改前**：
+```python
+_is_maca_target = target_is_maca(determine_target(return_object=True))
+```
+
+**修改后**：
+```python
+try:
+ _is_maca_target = target_is_maca(determine_target(return_object=True))
+except Exception:
+ _is_maca_target = False
+```
+
+**作用**：模块导入时如果 maca target 不可用，静默降级为 `False`，避免阻塞整个包的导入。
+
+### 5.3 修复 2：determine_target() 降级保护
+
+**文件**：`tilelang/utils/target.py`
+
+**修改前**：
+
+```python
+if is_maca_available:
+ return_var = "maca"
+```
+
+**修改后**：
+
+```python
+if is_maca_available:
+ try:
+ Target("maca")
+ return_var = "maca"
+ except Exception:
+ is_maca_available = False
+if is_maca_available:
+ return_var = "maca"
+```
+
+**作用**：返回 "maca" 之前先验证 `Target("maca")` 是否真的能创建成功，不能则降级到 CUDA/HIP/Metal 等其他后端。
+
+---
+
+## 6. 启用 MACA 后端重新编译（最终成功）
+
+### 6.1 关键发现
+
+查看 `CMakeLists.txt` 和 `cmake/FindMACA.cmake`：
+
+```cmake
+# CMakeLists.txt
+elseif($ENV{USE_MACA})
+ set(USE_MACA ON)
+
+# cmake/FindMACA.cmake
+macro(find_maca use_maca)
+ if(IS_DIRECTORY /opt/maca)
+ set(__maca_sdk /opt/maca)
+ endif()
+ find_library(MACA_MACAMCC_LIBRARY mcruntime ${__maca_sdk}/lib)
+endmacro()
+```
+
+**结论**：需要显式设置 `USE_MACA=ON` 环境变量，CMake 才会启用 MACA 后端。
+
+### 6.2 编译命令
+
+```bash
+cd /app/tilelang-metax
+rm -rf build
+export USE_MACA=ON
+pip install -e . -v
+```
+
+### 6.3 编译过程
+
+1. **安装构建依赖**：`scikit-build-core`, `cython`, `z3-solver`, `patchelf`
+2. **CMake 配置**：
+ ```
+ -- Found MACA_INCLUDE_DIRS=/opt/maca/include
+ -- Found MACA_MACAMCC_LIBRARY=/opt/maca/lib/libmcruntime.so
+ ```
+3. **Ninja 编译**：692 个编译目标，约 1 分 38 秒完成
+4. **生成 wheel**：`tilelang-0.1.9+maca.gitf1ca0fb9-cp38-abi3-linux_x86_64.whl`
+
+### 6.4 版本变化
+
+| 阶段 | 版本号 | 说明 |
+|------|--------|------|
+| 未启用 MACA | `0.1.9+cuda.gitf1ca0fb9` | 无 MACA 后端 |
+| 启用 MACA 后 | `0.1.9+maca.gitf1ca0fb9` | ✅ MACA 后端已编译 |
+
+---
+
+## 7. 验证安装
+
+### 7.1 基础导入验证
+
+```bash
+python -c "import tilelang; print(tilelang.__version__)"
+# 输出：0.1.9+maca.gitf1ca0fb9
+```
+
+### 7.2 MACA Target 验证
+
+```bash
+python -c "from tvm.target import Target; t = Target('maca'); print(t)"
+# 输出：
+# maca -keys=maca,gpu -max_local_memory_per_block=4095 -max_num_threads=1024
+# -max_shared_memory_per_block=65536 -max_threads_per_block=1024
+# -mcpu=xcore1000 -mtriple=mxc-metax-macahca -thread_warp_size=64
+```
+
+### 7.3 Kernel 测试验证
+
+```bash
+cd /app/tilelang-metax
+python -m pytest testing/python/kernel/test_tilelang_kernel_gemm.py::test_gemm_f16f16f32_nn -v
+# 结果：PASSED ✅
+```
+
+---
+
+## 8. 测试通过率汇总
+
+| 测试模块 | 启用 MACA 前 | 启用 MACA 后 |
+|----------|-------------|-------------|
+| `testing/python/cpu/` | 14/14 ✅ | 14/14 ✅ |
+| `testing/python/arith/` | 87/87 ✅ | 87/87 ✅ |
+| `testing/python/kernel/` | ❌ 全部失败 | **5 passed, 2 failed, 9 skipped** |
+| `testing/python/transform/` | ❌ 8 errors | **164 passed** |
+| `testing/python/language/` | ❌ 141 failed | **330 passed** |
+
+**仅剩失败**：
+- `test_gemm_i8i8i32_nt` — INT8 精度偏差
+- `test_gemm_i8i8i32_tn` — INT8 精度偏差
+
+---
+
+## 9. 复现步骤（一键脚本）
+
+```bash
+#!/bin/bash
+set -e
+
+# 1. 进入源码目录
+cd /app/tilelang-metax
+
+# 2. 清理旧构建
+rm -rf build
+
+# 3. 启用 MACA 后端
+export USE_MACA=ON
+
+# 4. 编译安装（可编辑模式）
+pip install -e . -v
+
+# 5. 验证导入
+python -c "import tilelang; print('Version:', tilelang.__version__)"
+python -c "from tvm.target import Target; print('MACA target:', Target('maca'))"
+
+# 6. 运行核心测试
+python -m pytest testing/python/kernel/test_tilelang_kernel_gemm.py::test_gemm_f16f16f32_nn -v
+
+echo "TileLang MACA build completed successfully!"
+```
+
+---
+
+## 10. 关键踩坑记录
+
+| 坑 | 现象 | 根因 | 解决 |
+|----|------|------|------|
+| **Ninja vs Makefile 冲突** | `ninja: error: Makefile:5: expected '='` | build 目录残留旧 Makefile | `rm -rf build` |
+| **未启用 MACA** | `Target kind "maca" is not defined` | 缺少 `USE_MACA=ON` | 重新导出环境变量并编译 |
+| **导入时急切检测** | `import tilelang` 崩溃 | `gemm_sp.py` 模块级别求值 | `try-except` 保护 |
+| **auto-detect 降级缺失** | 所有测试崩溃 | `determine_target` 直接返回未注册的 target | 创建前验证 target 可用性 |
+
+---
+
+*文档由实际操作过程整理，基于 tilelang-metax dev 分支（commit f1ca0fb9）*
\ No newline at end of file
diff --git a/tilelang-llm-inference-operator-optimization/基于国产软件栈的推理前沿算子优化比赛方案.md b/tilelang-llm-inference-operator-optimization/基于国产软件栈的推理前沿算子优化比赛方案.md
new file mode 100644
index 0000000..c360d6b
--- /dev/null
+++ b/tilelang-llm-inference-operator-optimization/基于国产软件栈的推理前沿算子优化比赛方案.md
@@ -0,0 +1,225 @@
+# 题目编号：XH-202608
+
+# 基于国产软件栈大模型推理前沿算子优化
+
+## 比赛方案
+
+### 一、发榜单位
+
+沐曦集成电路（上海）股份有限公司
+
+### 二、题目名称
+
+基于国产软件栈大模型推理前沿算子优化
+
+### 三、题目介绍
+
+我国在AI算法创新与应用落地方面已取得显著成就，主流大模型在算法层面已与国际先进水平并跑乃至部分领跑。然而绝大多数可公开下载和部署的大模型，其底层推理代码与性能优化均深度适配于国际主流AI生态。这种深度绑定，在中美科技竞争长期化、复杂化的今天，已成为我国人工智能产业自主发展的潜在风险。在此背景下，需要构建基于全国产软件生态的核心算子库，为国产大模型和国产算力芯片的协同创新与规模化应用，浇筑自主可控的软件"地基"。
+
+在此背景下，国产自主研发并于2025年开源的TileLang，作为一种专为AI算子开发设计的领域专用语言（DSL），其核心是将高性能计算中的"分块技术"（Tile）作为首要优化对象，自动推导并实施适合目标硬件的优化策略，让开发者更专注于算法逻辑本身。TileLang具有高效率、自动化、强兼容和开放生态的特性，成为构建国产软件生态中重要的开源项目。沐曦股份MXMACA软件栈已与TileLang完成深度适配和验证，可作为后端接入TileLang，构建贯穿上层框架到底层硬件的、全国产化的AI软件生态。
+
+本次比赛聚焦于大模型推理中3个关键前沿算子，以TileLang作为核心开发语言，以沐曦股份MXMACA作为软件栈后端，以曦云系列C500处理器为硬件算力，开展基于全国产软件生态的深度性能优化与实践。选择的算子包括：
+
+1. Fused Moe Gemm
+2. DeepSeek V3/R1所用Multi-Head Latent Attention（QK dim576 + VO dim512）
+3. Native Sparse Attention
+
+其中Fused moe gemm是一种针对混合专家模型的关键核心算子。混合专家模型通过门控网络动态选择少数专家进行处理，通过稀疏激活特性带来参数规模的优势，但也带来计算细碎化，硬件算力资源利用不充分的问题。Fused moe gemm通过计算融合优化技术解决这一技术挑战，该算子输入多个独立的专家权重矩阵、经过路由分配的输入激活张量，输出所有专家各自的计算结果，经后续处理（如聚合）后形成该MoE层的最终输出。该算子的技术核心在于通过将多个独立的专家矩阵乘法（GEMM）分组并融合为一个统一的内核进行计算，提升硬件算力利用率，加速大模型推理。
+
+DeepSeek V3中的MLA（Multi-head Latent Attention）算法通过将KV cache压缩为低秩潜在向量，在保持与标准多头注意力相当性能的同时，将KV cache显存占用大幅减小，显著降低了长序列推理的显存瓶颈；其创新的解耦旋转位置编码和优化的矩阵吸收机制，不仅提升了推理效率，还通过减少访存次数和增强计算强度，在大模型场景下实现了更高的吞吐量和更低的延迟，特别适合需要处理超长上下文的应用场景。有赖于MLA算法和MOE等关键技术，DeepSeek V3-671B的性能在很长一段时间内领跑国际先进开、闭源模型的水平。
+
+DeepSeek提出的Native Sparse Attention算法通过硬件原生的稀疏计算模式，将超长文本的注意力计算复杂度从O(n²)降至O(n√n)，在保持模型性能几乎无损的同时，实现了数倍的推理速度提升和50%以上的显存节省；NSA创新地通过分块稀疏模式和动态稀疏度调整机制，不仅避免了传统稀疏注意力中的负载不均衡问题，还通过优化的内存访问模式显著提高了GPU利用率，特别适合长上下文场景下的高效推理。实验表明，在64k长度序列上，NSA在解码、前向传播和反向传播各个阶段都实现了显著加速，同时在通用基准测试、长上下文任务和指令推理方面保持或超越了全注意力模型的性能，有效解决了长上下文建模中的计算瓶颈问题。
+
+本次比赛聚焦于上述算子的性能优化，要求参赛选手理解算子的计算逻辑， TileLang的基本语法和GPU算子的常见优化方法，在保证精度的同时提升算子性能。大赛将提供编程环境，功能和性能测试集，助力团队快速上手并开展实践。
+
+### 四、参赛对象
+
+参赛对象为2026年6月1日以前正式注册的国内全日制非成人教育的普通高等学校在校专科生、本科生、硕士和博士研究生（不含在职研究生），以及全日制职业教育本科、高职高专在校学生，可通过学生赛道申报作品参赛。
+
+参赛对象可以团队或个人形式参赛，每个团队不超过10人，每件作品可由不超过3名指导教师进行指导。可以跨专业、跨学校、跨地域组队，但同一团队所有成员均应符合本赛道相关年龄、身份要求。每件作品只可由1所高等院校、科研院所等作为参赛主体提交申报。
+
+### 五、答题要求
+
+#### （一）初赛作品要求
+
+根据提供的基础样例，使用TileLang在MXMACA平台上实现和优化Fused Moe Gemm算子，并需要在基础样例Fused Moe Gemm算法测试API端到端耗时性能上有提升作为晋级决赛的基本要求。评判标准分两个层级：基本功能和性能加分层级。
+
+1. **基本功能层级评判标准**：
+ Fused Moe Gemm功能test集通过率99%以上，并有参考样例性能的提升。
+
+2. **性能加分层级评判标准**：
+ 在满足基本功能层级评判标准的基础上，全部参赛选手实现的Fused Moe Gemm算法测试API端到端耗时（warmup 10次，重复100次取平均耗时），按耗时长短排名，耗时越短排名越高，并且会在决赛有10分评分会根据初赛性能排序作为依据如下：
+ - 第1名 +10分
+ - 第2名 +9分
+ - 第3名 +8分
+ - 第4名 +7分
+ - 第5名 +6分
+ - 第6名 +5分
+ - 第7名 +4分
+ - 第8名 +3分
+ - 第9名 +2分
+ - 第10名 +1分
+
+#### （二）决赛作品要求
+
+1. 根据基础参考样例，使用TileLang提交MLA或者NSA算子实现并优化性能，考虑到2个算子难度有差异，请选手从2个算子任选一个作为决赛参赛评比（注意：MLA技术分最高80分，则决赛最终80*0.6=48分），2个算子评价如下：
+
+ **使用TileLang在MXMACA平台上实现和优化DeepSeekV3 MLA算子【技术分满分：80分】**，技术分拆分如下：
+
+ - **基本功能层级评判标准（技术分：40分）**：
+ MLA功能test集通过率99%（按照通过case数目折算得分，例如：80%覆盖，40*80%=32分）。
+
+ - **性能加分层级评判标准（技术分：40分）**：
+ 在满足基本功能层级评判标准的基础上，全部参赛选手实现的MLA算法测试API端到端耗时（warmup 10次，重复100次取平均耗时），按耗时长短排名，耗时越短排名越高：
+ - 第1名 +40分
+ - 第2名 +38分
+ - 第3名 +36分
+ - 第4名 +34分
+ - 第5名 +32分
+ - 第6名 +30分
+ - 第7名 +28分
+ - 第8名 +26分
+ - 第9名 +24分
+ - 第10名 +22分
+ - 第11名 +20分
+ - 第12名 +18分
+ - 第13名 +16分
+ - 第14名 +14分
+ - 第15名 +12分
+ - 第16名 +10分
+ - 第17名 +8分
+ - 第18名 +6分
+ - 第19名 +4分
+ - 第20名 +2分
+
+2. **使用TileLang在MXMACA平台上实现和优化NSA【技术分满分：100分】**，技术分拆分如下：
+
+ - **基本功能层级评判标准（技术分：40分）**：
+ NSA forward API功能test集通过率99%（按照通过case数目折算得分，例如：80%覆盖，40*80%=32分）。
+
+ - **性能加分层级评判标准（技术分：60分，例如：满分客观分为：50*0.3=15分）**：
+ 在满足基本功能层级评判标准的基础上，全部参赛选手实现的NSA算法使用64K seqlen的输入测试API端到端耗时（warmup 10次，重复50次取平均耗时），按耗时长短排名，耗时越短排名越高：
+ - 第1名 +60分
+ - 第2名 +57分
+ - 第3名 +54分
+ - 第4名 +51分
+ - 第5名 +48分
+ - 第6名 +45分
+ - 第7名 +42分
+ - 第8名 +39分
+ - 第9名 +36分
+ - 第10名 +33分
+ - 第11名 +30分
+ - 第12名 +27分
+ - 第13名 +24分
+ - 第14名 +21分
+ - 第15名 +18分
+ - 第16名 +15分
+ - 第17名 +12分
+ - 第18名 +9分
+ - 第19名 +6分
+ - 第20名 +3分
+
+### 六、作品评选标准
+
+提交作品最终评价及权重如下：
+
+| 类别 | 评审维度 | 权重 | 说明 |
+|------|----------|------|------|
+| 客观评测 | 技术实现分 | 60% | 根据决赛选择的题目的技术得分*0.6给出技术实现分（备注:MLA技术最高分80分，NSA技术分最高100分）。 |
+| | 初赛排名加分 | 10% | 根据初赛名次取前10规则给与相应分数。 |
+| 主观评测 | 展示质量 | 20% | 根据提交代码质量、Demo 视频与测试报告完整度给分。 |
+| | 创新性与可扩展性 | 10% | 是否体现独特机制、具备推广潜力和上游合并社区价值。 |
+
+### 七、作品提交时间
+
+2026年5月至9月上旬，各参赛团队选择榜单中的题目开展研发攻关，各高校、科研机构等组织协调机构应组织学生和青年科技工作者参赛，安排专业人员给予指导，为参赛团队提供支持保障。
+
+2026年9月，各参赛团队要向发榜单位完成作品提交，具体要求详见作品提交方式。
+
+2026年9月，由发榜单位完成初审，确定入围终审擂台赛的晋级作品和团队。
+
+2026年10月，安排专门团队提供帮助和指导，各晋级团队完善作品。
+
+2026年11月，组织终审擂台赛，角逐"擂主"。
+
+### 八、参赛报名及作品提交方式
+
+#### （一）报名方式
+
+1. 参赛选手登录"挑战杯"官网www.tiaozhanbei.net，在"揭榜挂帅"擂台赛报名入口注册账号，登录大赛申报系统在线填写报名信息。报名信息提交后，下载打印系统生成的报名表。
+2. 申报人在报名表对应位置加盖所在学校或所在单位公章。
+3. 将盖章版报名表扫描件上传至报名系统，等待系统审核。请参赛选手注意查看审核状态，如审核不通过，需重新提交。
+4. 系统开放报名时间为2026年5月30日—6月30日，逾期后系统将自动关闭报名功能。
+
+#### （二）作品提交方式
+
+参赛团队应于规定时间内将申报作品统一打包压缩提交至大赛申报系统以及"挑战杯"官网。作品的提交除提到的客观评测外，在初赛/总决赛截止时间点前，参赛团队应将所有要求的材料，包括技术方案文档 (PDF)、日志与结果文件、展示视频（10分钟，总决赛）、算法介绍 PPT（总决赛）、源代码或 Notebook（可选，建议提交或开源）打包成 .zip 格式压缩包提交至 opensource@metax-tech.com。压缩包命名方式为：申报人所在单位－申报人姓名－作品名称－联系电话（例如：XX 大学－张XX－XX方案－手机号)。
+
+如出现作品相关材料因"挑战杯"官网附件大小容量限制无法提交的情况，可提交一份无法提交的情况说明，附无法提交界面截图，盖校团委章后扫描程pdf，作为作品材料提交至"挑战杯"官网，同步将作品相关材料全部发送至上述制定邮箱。
+
+1. 算法设计报告：详细说明kernel设计的技术方案、创新点和实现步骤，给出算法伪代码
+2. 功能测试报告：针对发榜单位提供的功能测试集合，提供完整的参数设置、计算结果的精度与golden校验对比报告
+3. 性能测试报告：针对发榜单位提供的性能测试集合，提供完整的参数设置、性能测试结果分析报告
+4. 源代码：用TileLang实现算子的完整源码代码（包括功能和性能测试源码），编译、运行test/benchmark方式的README.md文档
+
+除参赛报名表外，各参赛组提交的文档、源代码和模型文件不得携带任何参赛学校、老师和学生的个人信息。同时，各参赛团队在提交作品时，同步报送1份经报名系统审核通过的参赛报名表，报名表所有信息须与系统内填报内容完全一致。
+
+### 九、赛事保障
+
+1. **算力资源支持**：参赛选手在完成报名后，提供对应线上的曦云C500在线算力资源。
+2. **技术培训**：赛前赛中至少组织 2 场线上培训，提供培训回放及答疑文档；另外届时会根据实际情况，决定是否组织线下的专场培训。
+3. **专家指导**：建立线上答疑社群，由沐曦股份技术团队指导，定期回复技术问题。
+4. **交流平台**：在沐曦股份开发者社区开设赛事专属社区板块，支持参赛团队分享经验、交流问题，促进技术共创。
+5. **技术文档和课程**
+ 1. 提供 TileLang /MXMACA 技术文档
+ 2. 提供算子 baseline 示例
+ 3. 提供测试集、评测脚本
+ 4. 配备技术答疑团队
+ 5. 组织线上技术讲解/ Q&A
+
+### 十、设奖情况及奖励措施
+
+#### 1. 设奖情况
+
+按参赛作品数量比例设奖，原则上评出 "擂主"1个、特等奖5个，一等奖5个、二等奖6个、三等奖8个，获奖比例不超过参赛作品总数的30%，从特等奖中角逐出擂主团队。
+
+#### 2. 奖励措施
+
+- **奖金奖励**：擂主奖励10万元/个（叠加特等奖后奖金），特等奖奖励2万/个，一等奖奖励1万元/个，二等奖奖励0.5万元/个，三等奖奖励0.2万元/个。
+- **高端GPU奖励**：擂主团队2张GPU加速卡（不叠加特等奖GPU奖励），特等奖团队1张GPU加速卡。
+- **实习奖励**：优秀获奖者有机会参与之江&沐曦股份 "南湖之新"联合培养计划；所有获奖团队可获得赛事荣誉证书。
+- **曝光支持**：获奖作品可在沐曦股份开发者社区展示，作品可提供成果孵化与应用推广支持。
+
+备注：从特等奖中角逐出擂主团队，奖金10万元已叠加特等奖奖金，擂主不叠加特等奖GPU奖励。
+
+#### 3. 奖金发放方式
+
+比赛结束后，单位比赛专班工作人员与获奖团队取得联系，填写奖金申请表，赛事终审结果公示无异议后（公示期约1个月），在30个工作日内通过银行转账一次性发放至团队负责人指定账户，上述所列奖金均为税后奖金。
+
+### 十一、比赛专班联系方式
+
+#### 1. 专家指导团队
+
+- 顾问专家：马老师，联系电话：13811784391
+- 顾问专家：芦老师，联系电话：15395820130
+- 顾问专家：刘老师，联系电话：18210506627
+
+负责比赛期间技术指导保障。
+
+#### 2. 赛事服务团队
+
+- 联络专员：杨老师，联系电话：15201842467
+- 联络专员：章老师，联系电话：13501701786
+
+负责比赛期间组织服务及后期相关赛务协调联络。
+
+#### 3. 联系时间
+
+比赛期间工作日（9:00－17:00）
+
+### 附：发榜单位简介
+
+沐曦集成电路（上海）股份有限公司（股票代码：688802.SH）成立于2020年9月，于2025年12月成功登陆科创板。总部位于上海，并在北京、南京、成都、杭州、深圳、武汉、长沙等地设立全资子公司及研发中心。作为一家专注于全栈GPU芯片及解决方案的集成电路设计企业，致力于打造世界一流的GPU芯片及计算平台，成为数字经济的算力基石。
+
+公司拥有技术完备、设计和产业化经验丰富的团队，核心成员平均拥有近20年高性能GPU产品端到端研发经验，曾主导过十多款世界主流高性能GPU产品研发及量产。目前，公司已推出全面覆盖人工智能训练和推理、通用计算、图形渲染和科学智能等场景的四大序列产品，并配套自研 MXMACA 软件栈，真正实现了"软硬协同"，满足"高能效"和"高通用性"的算力需求。2025年起，公司坚持"开放协同、自主可控"的方向，全面推进计算生态以及产业生态的建设。在计算生态层面，公司以自主研发的 MXMACA 全栈软件栈为核心，积极拥抱开源，打造自主、开放、兼容的通用计算开源生态；在产业生态层面，公司依托"1+6+X"战略布局，以数字算力底座为基，持续深耕金融、医疗健康、能源、教科研、交通、大文娱六大重点行业，同步积极探索具身智能、低空经济等新兴领域，为数字经济与新质生产力发展提供坚实算力支撑。