从 Triton 编译链路切到 HIP 后出现核函数加载错误，通常不是 HIP 本身无法编译，而是两个编译体系对内存布局、栅栏语义和运行时初始化的假设不一致。先放下“换成 HIP 就能绕过 Triton 问题”的预期，按“编译产物是否生成、符号是否对齐、启动方式是否匹配”三步排查，大部分加载错误都能定位到具体环节。
切换 HIP 后核函数加载错误，多半源于编译产物与运行时环境不匹配，常见三类原因：未正确链接 HIP 运行时、核函数符号被裁剪或命名修饰不一致、显存分配或 stream 上下文与 Triton 残留状态冲突。建议先检查编译日志中的符号导出，再用最小可运行内核独立复现，避免直接套用 Triton 的启动模式。
先确认编译失败发生在哪一层
“Triton 编译算子失败”和“切换 HIP 后加载错误”是两个不同阶段的问题。Triton 失败可能发生在前端解析、中间表示生成或代码生成阶段，切到 HIP 后加载错误则发生在 .so 或 .hsaco 被加载到当前进程的时候。建议先区分：是编译阶段根本没有产出 HIP 源文件，还是编译出了文件但运行时加载报错。很多情况下，Triton 失败的真正原因是算子形状特化过多或使用了不支持的原子操作，切到 HIP 后依然要面对相同的算法逻辑，加载错误只是表象。
按三个方向排查核函数加载错误
1. 检查 HIP 运行时是否被正确链接和初始化
使用 hipcc 编译时，通常会自动链接 amdhip64。但如果工程中混用了 nvcc 风格编译选项，或手动指定了 -c 再单独链接，就容易漏掉 HIP 运行时。加载错误如果表现为“未定义符号”或“无法找到 amdhip64_*.so”，优先检查链接行。建议先用最简单的单文件内核，编译命令只保留 hipcc
验证方式：在运行环境中执行 ldd 或等价的依赖检查命令，确认可执行文件或动态库确实依赖 HIP 运行时。如果是动态加载 .so，还要确认加载时用的 RPATH 或 LD_LIBRARY_PATH 覆盖到 HIP 安装路径。
2. 检查核函数符号是否被导出并匹配调用侧预期
Triton 生成的核函数通常有固定的命名修饰规则，切换到 HIP 手写或半自动转换后，函数名可能被 C++ 修饰或宏改变。加载错误如果指向某个具体符号，用 nm 或 hipcc -### 查看对象文件中的符号，确认调用侧使用的名称和编译产物中的名称完全一致。注意 HIP 默认使用 C++ 链接，如果核函数是纯 C 接口，需要加 extern "C"。
特别是用 hipify 把 CUDA 代码转过来时，__global__ 或 __device__ 宏可能没有被正确替换，导致核函数被编译成了普通 C++ 函数，运行时找不到可启动的入口。这种情况下加载不会直接失败，但在 hipLaunchKernelGGL 调用时很容易报“invalid device function”一类错误。建议先打开编译器的 -Wall -Werror，把未使用函数或隐式声明这类警告当作错误处理。
3. 检查显存分配、stream 和模块加载的上下文是否干净
Triton 运行时会管理自己的缓存和显存池，切换 HIP 后如果进程还在，之前分配的 Triton workspace 可能和 HIP 的默认堆冲突。加载错误如果伴随非法地址或内存分配失败，可以尝试在加载 HIP 模块前先调用 hipDeviceSynchronize 或 hipCtxReset，清掉 Triton 留下的未完成操作。注意 hipCtxReset 会释放当前上下文的所有资源，如果其他模块正在使用，不能随意调用。
如果 HIP 核函数是从文件或 buffer 加载 .hsaco，运行时使用 hipModuleLoad / hipModuleLaunchKernel，需要确认 .hsaco 对应的 target 与当前 GPU 的 gfx 版本匹配。不匹配时会出现“no binary for current device”或类似错误。建议先运行 rocminfo 查看设备 ID，再在编译时指定正确的 `--offload-arch`，例如 `--offload-arch`=gfx90a 或 gfx1030，不要用默认的 gfx000。
一个最小验证序列
为了快速隔离问题，建议把验证拆成三步，每一步都能独立判断。
- 用 hipcc 编译一个简单向量加内核，使用 hipLaunchKernelGGL 启动，确认 HIP 基础链路可用。
- 将疑似失败的算子单独抽出来，保持相同的数据形状和数据类型，编译成独立动态库或可执行文件，确认加载错误是否仍然存在。
- 回到原工程，注释掉 Triton 初始化代码，只保留 HIP 相关路径，观察错误是否消失。
执行这三步时，记录每一步的编译命令和运行日志，不要只在原工程里加打印。比较失败模式的变化点，通常能锁定是编译选项、链接顺序还是运行时上下文的问题。
调试时的有效工具和环境变量
可以使用 hipcc -### 查看实际调用的子命令，确认编译参数是否完整。运行时设置 AMD_LOG_LEVEL=1 或 HIP_LAUNCH_BLOCKING=1 能强制内核同步执行，把异步错误变成同步错误，这样堆栈位置更容易对应到具体调用行。如果加载错误发生在模块加载阶段，还可以设置 HIP_MODULE_LOAD_VERBOSE=1（需要结合环境确认是否支持），查看模块加载器报出的具体阶段。
使用 rocprof 做性能或事件追踪不是排查加载错误的首选，但如果你怀疑错误是某个内核执行中途退出导致的，用 rocprof 或 hipMemset 加 hipDeviceSynchronize 可以提前捕捉到错误代码，避免错误被后续操作覆盖。
常见问题
切换到 HIP 后，核函数加载成功但启动时报“invalid device function”，怎么回事？
最常见原因是编译目标和当前 GPU 架构不匹配，或核函数实际没有被编译进模块。优先检查 hsaco 的 target 是否与 rocminfo 输出一致，再看核函数是否被内联或优化掉了。可以在函数前加 __attribute__((visibility("default"))) 或检查编译产物中的符号表。
Triton 和 HIP 可以在同一个进程中混用吗？
可以，但需要各自管理自己的设备状态。Triton 的缓存和 HIP 的上下文可能共享同一块显存池，建议在切换前显式同步设备。如果出现显存地址冲突，优先调整分配策略，而不是盲目重置上下文。