——AKO4X项目从CUDA到HIP的全栈迁移报告 模型：Claude Code + DeepSeek V4-pro | 环境：海光DCU K500SM\_AI (gfx928) + HIP # 一、项目背景与目标 AKO4X（Agentic Kernel Optimization — advanced & eXtensible）是一个面向高性能GPU算子的智能体驱动优化框架。其核心理念是将传统的人工kernel优化迭代流程——编写、性能分析、修改、再分析——进行结构化与自动化，通过引入编码智能体（如Claude Code）替代人工完成kernel的反复迭代优化。 当前，海光DCU生态相较NVIDIA CUDA生态仍处于早期发展阶段，算子开发与优化高度依赖工程师的手工迭代。开发者需要在编写kernel、运行hipprof性能剖析、分析瓶颈、修改代码之间反复循环，单个算子的优化周期往往需要数天甚至数周。将AKO4X迁移至DCU平台后，编码智能体可在结构化的闭环流程中自主完成这一迭代过程，显著缩短单算子的优化周期，降低对高级kernel工程师的人力依赖，并为在DCU平台上高效部署DeepSeek V4等大模型的推理服务提供底层算子基础。 # 二、CUDA到HIP迁移的核心挑战 AKO4X原本是为NVIDIA B200（Blackwell架构）设计的，整个项目不包含任何HIP编程相关的内容。迁移的第一步是系统性地梳理两个平台之间的不兼容性。经过分析，我们将差异归纳为三个层次：编程模型层面、硬件微架构层面和生态工具链层面。 ## 2.1 编程模型层面 HIP与CUDA的API转换是迁移中最直观的部分。HIP本身的设计即为CUDA的"方言"，大量API存在一一对应关系：cudaMalloc对应hipMalloc，cudaMemcpy对应hipMemcpy，cudaDeviceSynchronize对应hipDeviceSynchronize等。NVIDIA提供的hipify工具能覆盖大部分简单场景的自动转换。然而，"能编译通过"与"能正确运行"之间仍存在显著差距，真正的难点在于硬件层。 ## 2.2 硬件微架构差异——迁移的核心难点 **Wavefront与Warp的差异。**这是最基础也最关键的差异。NVIDIA GPU的最小调度单位是32线程的Warp，而海光DCU的是64线程的Wavefront。这意味着所有硬编码32的位置都必须调整为64，包括但不限于：Warp-level reduction中的步长（\_\_shfl\_xor\_sync的迭代次数从5次变为6次）、lane mask（从0xFFFFFFFF变为64位掩码）、以及Shared Memory的bank冲突分析模型。这并非简单的搜索替换所能解决——许多CUDA kernel在做Warp-level reduction时，会基于Warp size=32的假设来决定Shared Memory的分配大小、线程分组方式乃至循环边界。这些假设一旦移植到Wavefront=64的环境中，轻则导致Occupancy偏离、性能下降，重则引发reduction步数不完整的计算错误。 **Tensor** **Core差异。**NVIDIA Hopper架构的Tensor Core支持FP8（E4M3/E5M2）、BF16、FP16等多种精度的矩阵乘法，并通过mma.sync PTX指令和CUTLASS库提供了成熟的编程抽象。而海光DCU（gfx926/gfx928）的Tensor Core仅支持BF16精度，没有硬件FP8支持。在DCU上，Tensor Core的编程接口为DUMMA（DCU Matrix Multiply-Accumulate），其API形态与NVIDIA的WMMA/mma.sync完全不同。 **FP8格式的隐蔽差异。**即使DCU未来支持硬件FP8，也无法直接使用——AMD/CDNA系的原生FP8格式为FNUZ（Finite NaN, Unsigned Zero），而NVIDIA使用的是IEEE 754子集格式。两者的exponent bias和特殊值编码不同，同一个uint8字节在两种格式下解码出的浮点数值不一致。因此所有FP8相关场景在DCU上都需要走uint8 + 查表手动反量化的路径。 ## 2.3 一条已验证的死路 迁移之初，我们曾尝试通过DCU上的GPUfusion（NVIDIA兼容层）直接编译FlashMLA的CUDA源码。实测结果表明此路径完全不可行：FlashMLA大量使用了mma.sync等PTX级别内联指令以及CUTLASS的Hopper专用WGMMA/TMA原语，这些均超出兼容层的覆盖范围，编译直接报错且无workaround。这一结论决定了整个迁移的基本策略：不能"翻译"，必须"重写"。 # 三、迁移策略与架构设计 确认"必须重写"之后，我们采用了分层识别、正确性优先的策略来确定迁移范围和实施顺序。 ## 3.1 四层架构的迁移评估 **编排层（spawn.py）：**纯Python实现，不涉及GPU代码。迁移工作量较小，主要包括将GPU检测逻辑从nvidia-smi替换为rocm-smi，以及调整编译参数的默认值。 **模板与知识库（templates/）：**包含智能体的System Prompt和SKILL Catalog。原版覆盖Triton、CUDA、CuTe DSL等NVIDIA生态工具。迁移的核心工作是编写DCU专用的SKILL文件，涵盖HIP kernel开发规范、DUMMA Tensor Core编程、LDS优化策略及hipprof性能分析方法。这部分属于知识工程而非代码改动，但SKILL的质量直接影响智能体的优化效果。 **脚本层（scripts/）：**运行时工具，核心是benchmark\_adapter.py。需要将benchmark后端从flashinfer替换为hipBLAS或自建的HIP benchmark harness，同时将NCU profiling脚本全部替换为hipprof/hipprof\_ctrl的调用方式。 **Kernel代码层：**这是工作量最大的部分，也是迁移的核心战场。 ## 3.2 Kernel层迁移——以GEMM为例 以GEMM kernel为例，原版代码使用CUDA + CUTLASS/CuTe编写，深度绑定了Hopper的WGMMA和TMA。迁移到DCU上需要完成以下关键改造：将WGMMA替换为DUMMA，调整Tile尺寸以适配DCU的Tensor Core规格；根据DCU的LDS大小和bank结构重新计算TM/TN/TK参数和LDS buffer大小；根据DCU的内存层级重新调优全局内存的向量化加载宽度、对齐要求和coalescing策略。 这部分工作我们实际使用了Claude Code进行智能体驱动的迭代优化——大约100分钟的自动化调优过程中，智能体经历了约23个kernel版本的迭代。 # 四、迭代优化的经验总结 在GEMM kernel的DCU优化过程中，我们积累了一系列关于"什么有效"和"什么是死路"的工程经验。 ## 4.1 有效的优化手段 **向量化全局加载（uint4宽度）**带来了约3倍的性能提升，是单一最大的优化来源。合理选择TM/TN/TK的Tile尺寸，使LDS用量控制在硬件限制范围内、同时最大化Occupancy，也是关键因素。 ## 4.2 已验证无效的方向 **双缓冲LDS（Ping-pong Buffering）：**在当前矩阵尺寸和DCU配置下，双缓冲反而导致性能下降。原因是LDS占用翻倍导致Occupancy大幅下降，抵消了流水线化带来的延迟隐藏收益。 **TK=32：**将K维Tile从16调整到32，性能反而下降。根本原因是更大的TK增加了寄存器压力，导致寄存器溢出到Scratch Memory。 **sB的+1 Padding：**该方案试图消除bank冲突，但不仅未能提升性能，反而导致性能暴跌——因为padding打破了uint4向量化加载的对齐。 当前最佳结果达到约32.5 TFLOPS，约为hipBLAS的0.50倍。考虑到hipBLAS是厂商使用汇编级别手工调优的极致实现，且我们的kernel仍有prefetch、software pipeline、精细寄存器分配等优化空间，这是一个合理的起点。 ## 4.3 经验传承机制 我们建立了一套制度化的经验传承机制：每轮优化结束后，由智能体生成一份"经验沉淀文件"（ITERATION\_LOG.md）。该文件并非给人阅读的总结，而是写给"下一个接手优化的智能体"的交接文档。其核心要求包括：从真实代码和git log提取参数而非信任文字总结；重点记录每个失败方向的具体改动、结果和失败原因；钉死可复现状态（commit hash、确切参数、逐workload的TFLOPS数据）；约定追加而非覆盖以保留完整历史。 这套机制的价值在于：智能体驱动的优化本质上是一种搜索过程——大量尝试、快速验证、保留最优。如果每次新session都从零开始搜索，效率会很低。经验文件作为搜索空间的剪枝依据，可以显著提升后续迭代的收敛速度。 # 五、初步性能评估 目前已完成RMSNorm、GEMM和DSA Sparse Attention三个算子的迁移测试，所有算子的性能基线均为PyTorch。 **RMSNorm：**融合后的RMSNorm算子性能达到ATen官方融合kernel的3.4倍。需要指出的是，经过带宽测试分析，硬件利用率并不算高，性能优势更多来源于AKO4X生成的代码比ATen默认实现更适配gfx928架构特性。 如下图所示，hip为AKO4X编写的算子，naive为Aten官方融合算子   **GEMM：**在大矩阵场景下，算子性能尚未超过PyTorch，约为hipBLAS的0.50倍。这一结果表明仍有较大的优化空间。  **DSA** ******Sparse** **Attention：**该算子较为复杂，目前在正确性验证阶段尚未通过。 # 六、当前不足与原因分析 当前结果与预期目标之间仍存在差距，经分析主要原因如下： **模型知识短板：**当前大语言模型对HIP编程的知识储备相对薄弱。现有SKILL文件中的内容由Sonnet 4.6模型提炼而成，缺少高级工程师凝练的优化精髓，导致智能体在优化过程中频繁走弯路或在关键决策点缺乏方向感。 **知识污染：**原项目中残留的NVIDIA GPU相关知识零散掺杂其中，对智能体的推理过程产生了干扰。 **搜索策略不成熟：**智能体在迭代优化时倾向于"逐个方法尝试"的盲目搜索策略，缺乏基于硬件模型的方向性判断。 # 七、后续优化路线 **SKILL分级体系：**将现有SKILL拆分为基础知识层和高级优化技巧层。基础层覆盖HIP编程规范和DCU架构特性；高级层包含针对特定算子类型的优化方法论和经过验证的最佳实践，为智能体提供更扎实的知识库。 **引入预测性优化流程：**在进入下一次迭代优化之前，要求智能体先基于硬件模型预判优化方向的效果。如果分析表明性能大概率下降，则主动规避；如果预期效果不确定或有提升可能，才进行实际尝试。以此降低盲目搜索的比例。 **引入参考实现蒸馏：**尝试将Composable Kernel（CK）或rocBLAS中的高性能kernel作为参考实现提供给智能体，并开启自进化功能，观察智能体能否从这些工业级实现中自主蒸馏出可迁移的优化技巧。 **本文介绍了AKO4X****智能体****驱动优化框架从NVIDIA** **CUDA****平台到海光****DCU****平台的迁移实践。**虽然当前性能结果距离生产可用仍有差距，但迁移过程中积累的工程经验——特别是关于DCU硬件特性的正面与负面知识、以及智能体经验传承机制的建立——为后续的持续优化奠定了坚实基础。随着SKILL体系的完善和优化流程的迭代，我们有信心逐步缩小与厂商手工调优实现之间的性能差距。