1. 先把概念捋清楚分离编译到底解决什么问题1.1 从“整包编译”到“分离编译”一次典型的项目迁移我在做基于HIP的异构计算项目时有一个很深的体会刚上手的时候大家几乎都习惯把所有kernel和设备函数丢在同一个翻译单元里用hipcc一把梭编译。前期的确很爽不用处理任何跨文件依赖改一个函数重新编一次整个文件就行。但项目一旦超过几千行或者团队里多个人同时改代码这种模式就会变得非常难受——每一次增量编译都要把整个翻译单元重新过一遍链接阶段还会因为设备代码全部内联生成出体积巨大的fat binary。这时候就需要分离编译。所谓分离编译对于HIP生态来说本质上就是把“编译”和“链接”两个阶段拆开让不同的编译单元各自生成设备对象文件然后在链接阶段再合并。这意味着你的kernel可以像普通C函数一样声明在头文件里定义在另一个.cpp里多个编译单元之间可以互相调用设备函数而不再受“所有设备代码必须写在一个文件里”的限制。我第一次把项目从单一翻译单元拆成模块化编译时编译时间从原来的1分40秒降到了25秒左右。这个收益在大项目里会随着文件数量增加而放大而且方便做增量构建不用每次改一行就全家重新编。1.2 谁在背后干活的HIP分离编译的工具链拼图要理解分离编译必须先认识负责这套流程的工具们。它们不是一个人而是一整条流水线。clang负责把.hip或.cpp文件里由__global__、__device__标记的代码编译成设备bitcode或cubin对象clang-offload-bundler负责把host对象和设备对象打包成一个文件clang-offload-binder在静态库场景里负责解包和重打包ld.lld是最终执行合并链接的武器它调用llvm-mc生成带GPU镜像的最终可执行文件。再加上hip-device-libs提供内置库比如hc_math、ocml相关的那一堆这整套东西缺一个都能让你在链接阶段欲哭无泪。很多人拿到hipcc的-c参数就以为是开启了分离编译其实不是。真正的分离编译意味着设备代码的链接不再由clang完成而是推迟到链接阶段由lld统一处理。这背后有一套offsload bundling机制在支撑待会儿我会展开讲。1.3 分离编译 vs 整包编译一张表看清差异对比维度常规整包编译分离编译Relocatable Device Code编译粒度单翻译单元多编译单元独立编译设备代码位置必须内联在同一文件可跨文件声明和调用增量编译效率低改一处全量重编高只重编受影响文件fat binary大小较大包含重复实例化较小链接期裁剪链接阶段复杂度简单编译器代劳较高需lld介入适合项目阶段小型、单人或原型中大型、多人协作、库封装实际项目里我通常建议个人小脚本或者跟CUDA代码一对一翻译的早期验证阶段可以先用整包编译把功能跑通。到了要出库、要打包、要给别人调用的阶段一定要切换到分离编译。2. 分离编译的完整流程从源码到可执行文件的三次变身2.1 编译阶段设备代码是怎样被单独抽出来的分离编译的第一步是给clang传-c并加上--offload-archgfx90a这类目标架构参数。这一步clang会同时做两件事编译host代码同时编译设备代码。host代码编译跟普通CPU代码没什么区别唯一需要注意的是它里面会出现一些外部符号这些符号指向一个跟__hip_fatbin相关的数据结构也就是设备二进制代码的容器。设备端的函数声明会被编译成外部引用等链接阶段再解析。设备代码则会被clang的AMDGPU后端转成LLVM bitcode再经过opt和llc处理成AMDGCN架构的对象文件。这里有一个有意思的细节分离编译模式下clang不会在编译期去解析设备函数的外部引用而是保留这些未解析的符号放进一个叫__hip_fatbin的section。这其实是把“链接”这个动作延后了。我自己在做这一步时很容易犯的错是少了-fgpu-rdc这个开关。HIP的分离编译跟CUDA的-rdctrue一个道理不加这个flag编译器会默认把所有设备代码视为静态内联遇到跨编译单元的设备函数调用直接报“无法解析的设备符号”错误。所以这个flag必须跟-c一起出现提醒编译器走可重定位设备代码路径。2.2 打包阶段clang-offload-bundler怎么把两个世界装进一个文件编译完成后你会得到一个包含host对象和设备对象信息的中间产物但它们是分离的。下一步是用clang-offload-bundler把它们绑到一起。这里我打个比方host对象就像是快递包裹的纸箱设备对象是里面的货物。bundler干的事就是在纸箱上贴一张清单然后用特定的格式把货物和清单绑在一起让快递系统能识别。具体的命令大致是这样clang-offload-bundler --typeo --targetshost-x86_64-unknown-linux-gnu,hip-amdgcn-amd-amdhsa--gfx90a --inputsmain.host.o --inputsmain.device.o --outputsmain.fatbin -bundle执行完之后main.fatbin里就有了两个子文件的对象内容。注意这个fatbin只是一个容器它不是一个可执行文件真正的链接发生在下一阶段。这个打包过程看似简单但坑也不少。最常见的是--targets里指定的bundle id和实际编译时用的架构不一致。比如你用--offload-archgfx906编译设备代码打包时却写成了gfx90a那么即使命令没有立刻报错等到链接阶段也会出现找不到bundle的诡异错误。2.3 链接阶段ld.lld如何把碎块缝合成完整的GPU镜像链接阶段是整个分离编译的临门一脚。当你的项目里有多个编译单元时每个单元都会生成一个自己的fatbin。ld.lld需要读入所有这些fatbin把宿主代码里的外部引用逐一解析掉同时把设备代码的多个bitcode或对象文件合并成一个完整的AMDGCN镜像。如果你用的是“非rdc模式”lld的工作会简单很多只需从fatbin中把设备对象提取出来检查符号是否都在同一个翻译单元内然后生成可执行镜像。但在rdc模式下lld会把所有设备bitcode合并跑一轮完整的链接优化再生成最终镜像。这一轮优化包括函数内联、常量传播、死代码消除等常规链接期优化。这里有个容易忽略的细节最终生成的可执行文件并不只是把设备代码塞进去那么简单。lld需要把合并后的设备镜像写进ELF文件的特定section同时更新host侧的__hip_fatbin指针和元数据确保运行时HIP runtime能够正确定位和加载GPU代码。所以如果你看到链接成功但运行时找不到kernel八成是section分配或bundle元数据对不上而这个问题用readelf -S看section列表就能快速定位。3. 实操环节一个最小分离编译工程完整跑通3.1 环境准备关于“有没有预编译的LLVM”这个问题很多人在这一步卡住是因为不知道哪里去搞一套能用的LLVM和HIP工具链。自己从源码编LLVM不是不行但耗时太长一编就是两三个小时而且很容易编出来的版本跟ROCm的runtime库不匹配。我的建议是优先使用ROCm官方发布的预编译工具链。AMD的ROCm安装包会自带一整套编译好的clang、lld、hipcc和HIP runtime版本对应关系也已经验证过。比如ROCm 5.6自带的是LLVM 16ROCm 6.x对应LLVM 17或更高版本。如果你的应用场景需要在多个ROCm版本之间切换也可以直接用LLVM官方发布的二进制发行版但要注意两点第一必须确认这个LLVM版本带了AMDGPU后端和HIP支持第二需要手工把HIP device libraries路径指到对应的版本。我踩过的坑是这样的一开始图省事直接用了系统自带的老版本clang结果所有的__device__函数都被当作普通host函数编译链接的时候报了一大堆未定义引用。后来换了ROCm配的clang一切都正常了。一句话总结分离编译这活儿编译器版本敏感度非常高预编译工具链能直接满足要求的就尽量别自己造轮子。环境变量配置上我习惯这样做export ROCM_PATH/opt/rocm export HIP_PATH$ROCM_PATH/hip export PATH$ROCM_PATH/llvm/bin:$ROCM_PATH/bin:$PATH有了这一套clang、clang-offload-bundler、hipcc就都能直接用命令找到了。3.2 一个跨文件调用的最小样例我准备了一个极简的例子模拟真实项目里最常见的场景一个kernel声明在头文件里定义在单独的源文件中main函数在另一个文件里。首先是kernel.h#ifndef KERNEL_H #define KERNEL_H __global__ void vector_add(float* a, float* b, float* c, int n); #endif然后是kernel.cpp#include kernel.h __global__ void vector_add(float* a, float* b, float* c, int n) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx n) { c[idx] a[idx] b[idx]; } }最后是main.cpp#include iostream #include hip/hip_runtime.h #include kernel.h int main() { const int n 1024; float *a, *b, *c; hipMalloc(a, n * sizeof(float)); hipMalloc(b, n * sizeof(float)); hipMalloc(c, n * sizeof(float)); std::vectorfloat h_a(n, 1.0f), h_b(n, 2.0f), h_c(n); hipMemcpy(a, h_a.data(), n * sizeof(float), hipMemcpyHostToDevice); hipMemcpy(b, h_b.data(), n * sizeof(float), hipMemcpyHostToDevice); vector_add1, n(a, b, c, n); hipDeviceSynchronize(); hipMemcpy(h_c.data(), c, n * sizeof(float), hipMemcpyDeviceToHost); for (int i 0; i 5; i) { std::cout h_c[i] ; } std::cout std::endl; hipFree(a); hipFree(b); hipFree(c); return 0; }这里我把vector_add声明和定义分开就是为了演示分离编译的场景。如果只用一个翻译单元kernel的定义就在main前面编译是顺理成章的。现在分开之后就需要我们手动走一遍编译、打包、链接的完整流程。3.3 手把手走流程从编译到可执行文件第一步编译各翻译单元。这里的关键是必须加-fgpu-rdcclang -c main.cpp -o main.o -fgpu-rdc --offload-archgfx90a -I$ROCM_PATH/include clang -c kernel.cpp -o kernel.o -fgpu-rdc --offload-archgfx90a -I$ROCM_PATH/include你可以用file main.o检查一下输出会发现它既不是纯host对象也不是纯设备对象而是包含了fatbin的混合体。如果用readelf -S main.o能看到__hip_fatbin这个自定义section里面存放的就是打包后的设备代码。第二步链接。这里我直接用hipcc它会自动调用底层的lldhipcc main.o kernel.o -o vector_add -fgpu-rdc --offload-archgfx90a如果一切顺利这步不会打印任何输出。如果缺少-fgpu-rdc你会看到大量类似“undefined symbol: __device_stub__Z10vector_add”这样的错误这说明设备代码编译路径不对host侧的stub函数没有被正确定义。第三步验证。可以用rocminfo确认GPU型号然后直接运行./vector_add输出应该是1 2 3 4 5。3.4 CMake工程里的配置方式手动敲命令行适合理解原理真实项目里我还是建议用CMake来管理尤其是当文件数量多起来以后。这里给出一个能直接用的CMakeLists.txt片段cmake_minimum_required(VERSION 3.21) project(hip_separate_demo LANGUAGES CXX) set(CMAKE_CXX_COMPILER hipcc) set(CMAKE_CXX_STANDARD 17) set(CMAKE_CXX_STANDARD_REQUIRED ON) set(CMAKE_HIP_ARCHITECTURES gfx90a) add_executable(vector_add main.cpp kernel.cpp) set_target_properties(vector_add PROPERTIES CXX_STANDARD 17 CXX_STANDARD_REQUIRED ON ) target_include_directories(vector_add PRIVATE ${ROCM_PATH}/include) target_link_directories(vector_add PRIVATE ${ROCM_PATH}/lib)注意我特意把gfx90a写死。如果你的机器是其他GPU可以用rocm_agent_enumerator查一下或者直接设置成native让hipcc自动探测。写死的好处是避免团队里不同机器编译出不同架构的二进制坏处是分发时需要针对目标GPU重新编。4. 实操过程与核心环节实现4.1 用clang-offload-bundler手动拆包和打包有时候你需要单独处理fatbin比如从一个对象文件中提取设备代码做分析或者把多个设备对象重新打包。这时候就需要手动操作bundler工具了。这里我以拆分刚才生成的kernel.o为例看看它内部到底打包了什么东西。mkdir -p extract clang-offload-bundler --typeo --targetshost-x86_64-unknown-linux-gnu,hip-amdgcn-amd-amdhsa--gfx90a \ --inputskernel.o --outputsextract/host.o --outputsextract/device.o -unbundle执行完extract目录下会有两个文件。host.o是纯host侧对象里面权当没有设备代码只有stub函数和fatbin符号device.o是纯AMDGCN对象可以用llvm-objdump -d extract/device.o查看里面的GPU指令这里面就是真正在GPU上跑的玩艺。反过来如果你把两个对象单独编译后想合并成一个fatbin用我刚才在第二章提到的命令做一次-bundle就行。实际项目里这个过程常用于封装第三方库。比如你拿到一个只有.a静态库、没有源码的HIP库链接时发现设备符号无法解析就可以拆开.a找到里面的设备对象重新打包一份包含完整设备代码的fatbin。4.2 链接器脚本和section layout的观察我在排查链接问题的时候最常用的工具是readelf加llvm-objdump。一个正常的分离编译可执行文件ELF文件里至少应该有这几个关键section.texthost侧代码.hip_fatbin打包后的设备代码容器.debug_*如果开了调试选项的话可以用这个命令快速观察readelf -S vector_add | grep -E hip|text|rodata正常情况下你会看到类似这样的一行[16] .hip_fatbin PROGBITS 0000000000026000 026000 01e000 00 0 0 16这个section的大小就是设备代码的体积。如果这个section是0说明设备代码没有被打包进去运行时hipLaunchKernel就会找不到kernel入口。另一个值得留意的点是.note段里的AMDGPU属性信息比如amdgpu.metadata。这个metadata记录了kernel的参数布局、分组大小、架构特性等如果链接阶段这条信息丢失运行时就会报“unsupported ISA”或“invalid kernel descriptor”之类的错误。遇到这类问题优先检查是不是链接时漏掉了--offload-arch。4.3 常见的链接错误与排查思路先说最大概率遇到的一类错误undefined reference to __hip_fatbin。这个错误字面意思是host侧代码引用了一个叫__hip_fatbin的符号但链接器找不到它的定义。为什么会找不到最常见的原因是你把设备代码没有打包进最终的对象文件。比如你用clang -c编译但忘了加-fgpu-rdc那么设备代码就只存在于编译期的临时文件里没有被写进对象文件。这种情况下检查一下编译命令确认-fgpu-rdc和--offload-arch都加上重新编译即可。另一类高发错误是unsupported bundle ID或could not identify target。这种问题多半是打包和链接时的target字符串不一致。比如编译时用的是hip-amdgcn-amd-amdhsa--gfx90a但链接时clang自动探测到的架构是gfx906lld在fatbin里找不到匹配的bundle自然就报错了。还有一类错误跟bitcode版本有关表现是链接时提示bitcode version mismatch或者LLVM IR generated by incompatible version。这种通常发生在混用了不同版本的clang和lld时。分离编译对编译器版本一致性要求很高编译时用的clang和最终链接时用的lld必须是同一个版本家族否则生成的bitcode和对象格式就可能不兼容。4.4 动态链接与静态链接的选择HIP工程在链接阶段还有一个容易困扰新人的维度设备代码是静态打进可执行文件里还是作为独立动态库存在常规情况是静态链接所有设备代码都合并进最终ELF的.hip_fatbin。但如果你在做一个插件系统希望不同插件各自携带kernel那么就需要将设备代码编译成独立的.so并在运行时动态加载。HIP runtime是支持这种模式的前提是链接动态库的时候同样走分离编译流程让每个.so都包含自己的fatbin。我在一个实际项目里就遇到过这个问题主程序加载插件每个插件里都有若干个kernel如果全部静态链接成一个文件插件更新一次就得全量重发。改用动态链接后每个插件独立出.so主程序通过dlopen加载更新插件时只需要替换对应的.so。这个方案在PCIe环境下完全可行而且省去了主程序反复编译的麻烦。要注意的是动态链接模式下插件.so和主程序必须用同一套HIP runtime库否则会出现两个runtime实例抢GPU资源的问题。解决办法是每个.so都显式链接libamdhip64.so让动态链接器自动处理依赖。5. 用分离编译组织真实项目的几个建议5.1 工程结构上如何组织设备代码分离编译带来了自由也带来了责任。我的建议是把设备代码划分成三层结构。最底层是“设备库层”代码里只有__device__函数和基本数据结构不包含任何__global__函数这一层主要给上层调用中间层是“kernel封装层”每个__global__kernel单独成一个编译单元定义清楚输入输出和grid布局最上层是“host调度层”负责内存管理、流同步和kernel launch。这种分层最明显的好处是编译依赖关系非常清晰。设备库层变更不会影响kernel层的调用逻辑kernel层变更不会影响host调度层的代码。特别是多人协作时每个人只需要关心自己负责的那一层合并代码时的冲突概率会小很多。另外建议把公共的设备函数声明放到一个专门的device_utils.h头文件里并加上#pragma once保护。头文件里只放声明不放定义定义统一放在.cpp或.hip文件里这样才能真正利用分离编译的跨文件调用能力。5.2 编译期优化选项怎么取舍分离编译模式下设备代码的链接期优化会做一次全局的函数内联和常量传播。这意味着你可以放心使用内联程度高的写法而不用担心代码膨胀问题。不过要记得把-O3同时传给编译和链接阶段。只优化编译阶段不优化链接阶段生成的设备镜像质量会大打折扣。我测试过一个计算密集型的example只优化编译阶段时kernel运行时间大约是1.35ms优化了链接阶段后降到了0.82ms差距非常明显。还有一个值得开的是-ffast-math如果算法不涉及特别严格的IEEE浮点语义这个flag能显著提升GPU代码性能。但慎开因为它会改变浮点运算的舍入行为可能在数值稳定性上产生微妙影响。5.3 调试与release模式的分离编译差异调试模式下设备镜像不会做激进的优化符号信息也更完整。用rocgdb调试时需要确保编译时加了-g和-O0并且链接阶段没有被lld做重大优化。这里有个很多人不知道的点分离编译的调试符号比整包编译更容易跟随因为每个编译单元保留了独立的debug信息rocgdb可以精确地加载对应源码行和局部变量。我当时调试一个设备端内存越界时靠着rocgdb的info locals命令很快定位到是一个索引变量在极端情况下超出了数组边界。Release模式下我会把-g去掉加-DNDEBUG并把架构参数固定为实际部署机器的型号。注意release模式下链接失败的概率比debug模式高因为优化器会移除更多“看起来没用”的代码如果你的kernel有特殊导出需求记得用__attribute__((used))显式保留。5.4 静态库场景链接动态链接库时要注意的坑很多项目会把设备代码打包成静态库分发给下游比如libmydevice.a。这时候要特别小心不要让clang把设备代码当作host代码处理也不要让ar在打包时去掉fatbin段。一个经过验证的流程是先分别编译出每个模块的.o确认每个.o里都有__hip_fatbin段然后再用ar rcs libxxx.a x.o y.o z.o打包。注意ar默认不会动自定义段所以这步一般没问题。但在链接时如果库里的某些符号没有被引用链接器可能不会把它对应的设备代码拉进来导致最终可执行文件里缺少某个kernel。解决办法是加--whole-archive强制链接所有成员或者用-Wl,--unresolved-symbolsignore-all配合链接脚本手动控制。这个坑不常见一旦遇到往往要折腾半天所以提前意识到就好。6. 参考配置与常见问题速查表6.1 环境配置速查配置项推荐值说明ROCm版本5.6及以上推荐6.x分离编译的稳定性更好LLVM版本与ROCm配套检查clang --version确认CMAKE_HIP_ARCHITECTURES具体GPU型号或native列表或空格分隔编译flag-O3 -fgpu-rdc二者缺一不可链接flag-fgpu-rdc --offload-arch与编译保持一致还有个容易被忽略的点是hipcc的--gpu-architecture参数优先级。如果你既设置了--offload-arch又设置了--gpu-architecture编译器会优先用后者两者不一致时会产生大量warning甚至错乱。我一般只用一个参数来控制架构命令行上写--offload-archCMake里只设置CMAKE_HIP_ARCHITECTURES。6.2 错误信息速查表错误信息常见原因解决办法undefined reference to__hip_fatbin编译时未加-fgpu-rdc重新编译加-fgpu-rdcundefined reference to__device_stub__...设备代码没走HIP路径确认clang带HIP环境用hipcc替代clangunsupported bundle ID打包/链接架构不匹配统一架构字符串用readelf -S查看现有bundlebitcode version mismatchclang和lld版本不一致用同一ROCm套件版本could not identify target缺少--offload-arch编译和链接都显式指定GPU架构GPU memory access fault内核越界或空指针用rocgdb调试检查索引计算ld.lld: error: section .hip_fatbin is too large设备镜像过大检查是否有重复模板实例调整-fvisibility选项运行时找不到kernelfatbin没被正确链接用readelf确认__hip_fatbin段存在hipMemcpy返回hipErrorInvalidValue主机或设备端指针无效检查hipMalloc返回值确认拷贝字节数正确多个.so重复初始化runtime主程序和插件各带一份runtime显式链接libamdhip64.so避免静态runtime6.3 一点总结性质的实战心得或者说“最后再做一遍检查”我个人的习惯是每次做分离编译无论项目大小都会跑一遍“三查”流程编译阶段查一下警告有没有“unused device function”之类的提示链接后readelf查看__hip_fatbin段是否存在大小是否合理运行时用rocprof或hipEvent测一下kernel能否正常启动和耗时是否符合预期。这三步看起来很基础但能把绝大多数坑提前拦下来。分离编译相比整包编译确实会在配置和使用上多花一些学习成本。但当你项目规模上来之后这个成本会被增量编译省下的时间、代码组织的清晰度、以及多人协作时的便利快速抵消。现在ROCm工具链对分离编译的支持已经很成熟了比起早期各种版本错配问题现在的坑更多是“参数没传对”这种小问题。如果你遇到的环境问题排查了半天都找不到头绪第一步先确认所有工具链来自同一个ROCm/LLVM版本第二步统一架构字符串第三步检查每个.o里是否都有__hip_fatbin段。这三板斧基本上可以解决掉九成以上的分离编译问题。