LLVM HIP分离编译全攻略：原理、实操与踩坑记录
1. LLVM HIP分离编译：为什么需要它，以及怎么落地
做GPU异构计算开发的兄弟，应该都有过这种体验：项目一大了，每次改一行头文件，整个工程从头编一遍，几分钟起步，CPU风扇狂转，人在边上干着急。特别是用到HIP（Heterogeneous-Computing Interface for Portability）的时候，默认的编译流程会把每个.cpp文件当成一个完整的HIP编译单元，device端和host端一起优化，控制流分析、内联、指令调度全走一遍，算力全花在重复劳动上。我最早接触HIP分离编译，就是在接手一个多模块的视觉算法库时被逼出来的。那个库有几十个源文件，公共头文件还层层嵌套，最初用hipcc直接编译，一次全量构建要将近二十分钟，增量构建也经常因为一个头文件的变化引发大半工程的连锁重编。后来把编译流程拆开，device代码和host代码分开处理，中间产物缓存起来，才算把构建时间压下来，日常改bug的效率也完全不一样了。
这篇东西就围绕“LLVM HIP的分离编译”展开。我会先讲清楚HIP编译链路的底层逻辑，再完整走一遍分离编译的实操方案，包括依赖管理、中间产物生成、静态库归档和链接阶段的关键点，最后把我在实际工程中踩过的坑和排查思路一并整理出来。适合正在用HIP做正式项目、觉得全量编译扛不住的开发者，也适合刚接触HIP、想搞明白hipcc到底做了什么的朋友。
2. HIP编译链路的底层逻辑：为什么默认会慢，分离到底分离什么
2.1 HIP编程模型和GPU编译的特殊性
HIP是AMD主推的异构计算编程接口，目标很直接：给开发者一套类似CUDA的编程体验，同时能跑在AMD ROCm平台和NVIDIA CUDA平台（通过HIP-CUDA后端）上。你在代码里写的__global__函数、__device__函数，最终都要被编译成GPU能执行的二进制指令，这部分叫device code；而普通的入口函数、显存管理调用、循环调度逻辑，跑在CPU上，这部分叫host code。同一份代码里，两类逻辑经常混在一起，编译器必须把它们区分出来，分别生成目标代码，再由运行时把两者衔接起来。
这里的关键在于：哪怕你只写了__global__标记的一个核函数，编译整个文件时，编译器也必须在device pass里做完整的优化。LLVM会把这部分函数从host的优化流程里剥出来，单独走一遍-target为amdgcn-amd-amdhsa的代码生成，包括对GPU架构特有指令的选择和调度。这比普通CPU编译多出不少工作量。
2.2 hipcc默认做了什么事：为什么慢
hipcc本质上不是一个新的编译器，它是clang的一个包装脚本。对.cpp文件，它会自动加上一堆HIP相关的include路径，设置好ROCm的链接库，然后把整个文件交给clang，让clang在同一个编译单元内同时处理host和device两端。
默认模式下，每个.cpp文件里只要含HIP头文件，就算你只写了一句printf，也会被完整走一遍GPU的编译流程。这带来两个结果：一是单个文件编译时间比普通C++编译长；二是所有include的头文件、模板实例化、内联函数优化，在每个编译单元里都要重复做一遍。一个工程几十个文件，每个文件都重复做相同的头文件解析和公共代码优化，时间自然就上去了。性能好点的开发机还能忍，到了CI或者大型服务器上，每次提交都在烧时间。
2.3 分离编译到底分离了什么
分离编译的核心思路和CPU端C/C++工程没什么本质区别：把“编译各源文件”和“链接最终产物”拆开，每个源文件先编译成目标文件（.o）或中间字节码（.bc），最后再统一链接。但在HIP场景下，这个“分离”多了一层含义——不只按文件分离，还要按代码类型分离。
具体来说，一个HIP源文件在分离编译模式下会经历这些阶段：
- 预处理和语法分析，识别出host代码和device代码的边界；
- host部分按CPU目标（比如
x86_64）进行编译和优化，生成host目标文件； - device部分按GPU目标（比如
gfx90a）进行编译、优化和指令选择，生成GPU可执行码； - 用工具把这些内容合并或归档，供链接阶段使用。
这样做的直接收益是：当某个文件只改了一行CPU逻辑，GPU部分完全没动时，增量构建不需要重新优化device代码，huge的时间就省下来了。更妙的是，你可以把稳定的device代码先编译好归档，不同模块之间按需引用，整个工程的可扩展性大大提高。
3. 实操准备：环境搭建和工具链确认
3.1 ROCm和LLVM版本的选择
做分离编译，先得把工具链搞对。我这里用的是ROCm 5.7.x，配套的LLVM/Clang版本是16/17这条线。AMD官方对ROCm版本有绑定关系，不建议自行混搭——我见过有人在ROCm 5.4的环境里硬上ROCm 6.0的hipcc，结果链接阶段出现一堆符号找不到的玄学报错。
验证环境是否正确的命令：
如果hipcc --version能正常输出版本信息，说明基础环境没问题。再用一个最简测试确认GPU代码真能跑：
这里有个容易踩的坑：只有编译生效还不够，运行时的ROCm驱动、device library、libhiprtc等运行时组件版本也得对齐。如果驱动版本和ROCm版本不匹配，启动时往往报HSA_STATUS_ERROR_DRIVER_VERSION_MISMATCH或类似的HSA错误，代码明明编译过了，就是跑不起来。
3.2 没有合适预编译LLVM时怎么办
关于热搜里提到的“有没有预编译的llvm”这个问题，很多从CUDA转到HIP的开发者都有这个疑问。答案是：有，但要用对路子。AMD在ROCm的release包里会带一套经过完整配置、包含AMD GPU后端和各种target描述文件的预编译LLVM/Clang，位置一般在/opt/rocm/llvm/bin下。直接用系统包管理器装的LLVM（比如apt install llvm）默认可能没有AMD GPU的target支持，或者缺少ROCm的device library路径。
我自己更推荐直接用ROCm自带的LLVM，原因很简单：版本匹配、target齐全、device library能对得上。如果因为某些原因必须用脚下的预编译版本，那编完以后务必确认以下几点：
clang -print-targets输出里能看到amdgcn这一项；clang --hip-version能输出HIP版本且与ROCm一致；clang -print-resource-dir指向的resource目录里存在amdgcn相关的bitcode文件。
少一个都可能让你在分离编译时卡在设备代码生成阶段。如果你需要定制LLVM，比如想在老版本LLVM上增加某个新GPU架构的支持，那就只能源码编译了。常规操作流程是：clone LLVM仓库，切到对应release分支，用-DLLVM_TARGETS_TO_BUILD="AMDGPU;X86"这种参数配置目标平台，然后编出来。这一步速度和磁盘开销都比较大，配置不复杂的机器大约得半小时到一小时。
4. 分离编译实操：手把手跑通全流程
4.1 源文件的设计和编译单元划分
先给一个具体的示例工程结构。假设我们要做一个小的图像处理库，包含两个模块：一个负责颜色空间转换，一个负责卷积滤波。目录长这样：
两个核函数分别声明在头文件里，定义在各自的.cpp中。main.cpp负责调用。
这种文件划分看起来简单，却是分离编译能不能受益的关键。如果一个文件里塞了太多不相关的函数，改动任何一个都会触发这个文件所有device代码的重编译，分离就失去了意义。建议按“业务模块+稳定边界”的原则切分源文件，尽量让头文件保持小而稳定，少include大而全的公共头。
4.2 用生成依赖文件控制增量编译
增量编译的基石是精确的依赖关系。和Makefile里手写g++ -MM类似，hipcc可以生成依赖文件。
这里几个参数的作用分别是：
-M系列：让编译器在编译的同时输出当前源文件依赖的头文件列表，-MF指定输出位置；--offload-arch=gfx90a：明确指定GPU架构，生成对应的device代码；-D__HIP_PLATFORM_AMD__：告诉HIP头文件当前是AMD平台。如果不加，有些版本的HIP头文件会自动判断，但显式声明更保险。
生成出来的.d文件长这样：
Makefile或Ninja读取这些.d文件，就知道哪个头文件变化了，对应哪些目标文件要重新编译。这个环节看似基础，但很多新手做分离编译就是从这一步开始出错——依赖文件没配置对，导致改了头文件却不触发重编，最后出现“编译过了但行为不对”的诡异bug。
4.3 真正的分离编译：只编device代码，host代码走普通编译
前面提到，默认hipcc -c会把整个文件当作HIP编译单元处理。分离编译的精髓在于：如果你的源文件里host逻辑和device逻辑是分散的，可以利用LLVM对HIP code对象（bitcode）的处理能力，把device部分编译出独立的中间文件，host部分用普通clang编译。
一个更细化的流程是这样：
第一步，用hipcc把一个.cpp编译成包含GPU code的fat binary目标文件：
第二步，如果某个文件完全不含GPU代码，就可以绕开hipcc，用系统clang直接编：
注意，这里main.cpp还是需要include HIP头文件，才能调用hipDeviceSynchronize()这类API。但因为它没有__global__函数，没有必要走完整的device pass。直接用clang++编普通CPU代码即可，前提是你能正确找到HIP的include路径。
这样做完之后，链接阶段就要特别小心了。hipcc在链接时不只是调用clang，它还要用ld.lld对GPU code object做device link。如果GPU device代码和host代码已经分开生成了，链接时要用hipcc --hip-link来把device代码链接成最终可执行的fat binary。正规做法仍然是用hipcc做最终链接，它会自动处理-L/opt/rocm/lib -lamdhip64这类库参数。
4.4 中间产物：目标文件、bitcode、静态库
分离编译的好处是可以把中间产物保存下来。不同场景下，中间产物的格式有不同的用途：
- 目标文件
.o：最常用的格式。包含了host代码和用于最终链接的GPU code object元数据。多个.o直接hipcc链接成可执行文件。 - LLVM bitcode
.bc：适合再做一层跨模块优化。比如你希望对不同编译单元做link-time optimization（LTO），可以统一输出.bc再统一链接优化。 - 静态库
.a：把一组.o打包成库，工程依赖管理就方便多了。生成方法和CPU静态库一样，还是llvm-ar或ar。
静态库的打包方法：
链接时：
这里有个细节：同一个项目的所有.o在打包成.a之前，最好确保用的是同一组HIP编译参数，尤其是--offload-arch。如果两个.o一个编的是gfx906，一个编的是gfx90a，最终链接时虽然编译器可能会都保留，但二进制体积会增加，运行时也可能因为fat binary里的多个code object切换带来不必要的开销。这种问题实际开发中特别隐蔽，编译不报错，性能却莫名变差。
5. CMake集成：让分离编译落地到正式工程
5.1 用CMake组织分离编译项目的结构
直接用命令行编完一次，下一步当然是把流程固化到CMake里。AMD官方其实提供了/opt/rocm/share/rocm/cmake下的工具链支持，只要设置好CMAKE_HIP_COMPILER就能让CMake识别HIP语言。
一个可落地的CMakeLists.txt大致长这样：
把color_convert.cpp编译成静态库，再把main.cpp链接进来，CMake会自动完成依赖跟踪和增量构建。如果某个文件不需要GPU编译，可以在set_source_files_properties里关闭它的HIP处理，直接用CXX编译它。
5.2 让--offload-arch和GPU架构对齐
这块值得单独说。--offload-arch（CMake里叫CMAKE_HIP_ARCHITECTURES）的值取决于你的目标GPU。如果写错了，最常见的现象就是运行时“no kernel image is available for execution on the device”。这是HIP/CUDA里最经典的报错之一，本质上就是fat binary里没包含当前GPU架构对应的code object。
确定当前GPU架构的方式：
输出里能看到类似gfx90a、gfx1030的字符串。如果你的机器有多张不同型号的卡，可以一次指定多个架构，用空格或分号隔开：
这会生成包含两个code object的fat binary，体积会变大，但兼容性更好。部署环境GPU型号不统一时，这是个稳妥的做法。
6. 常见问题与排查技巧实录
6.1 “没有预编译的llvm”到底怎么解决
这个困扰不少人的问题，其实要分情况看。
如果你的目的是跑标准HIP程序，那ROCm自带的LLVM就是最省心的。/opt/rocm/llvm/bin里那套clang自带AMDGPU target支持，device library也齐全，直接用。
如果你的目的是研究LLVM本身的编译流程，想做自定义的pass或深度调优，那还是得源码编译。源码编译时注意opcode的默认配置，假如你没有为AMDGPU和X86两个target执行LLVM_TARGETS_TO_BUILD配置，最后编译出来的编译器可能根本没有GPU后端的代码生成能力，分离编译等于空谈。
如果仅仅是图省事想用系统包管理器装一个现成的LLVM，建议先确认版本和ROCm的兼容性。很多人在这一点上栽过跟头——装了个新版本的LLVM，结果无法处理AMDGPU target的bitcode流，编译时报一堆奇怪的附录错误。我的建议是：非必要别混装，ROCm自己带的编译器最合适生产环境。
6.2 链接时报“undefined reference to __hip_xxx”怎么办
这个报错典型的原因有两个：
一是链接库顺序问题。静态库的链接顺序是右侧依赖左侧，也就是说被依赖的库要放在依赖者的右边。如果写hipcc main.o -lcolor_convert -lamdhip64，而color_convert里的GPU kernel又依赖amdhip64运行时，那-lamdhip64必须在-lcolor_convert的右边。顺序反了，符号表解析不了，就会出现undefined reference。这其实是老生常谈的C/C++静态库链接顺序问题，到了GPU这里一样存在。
二是漏了--hip-link参数或缺少设备链接。手动分离编译时，如果只把.o简单打包，而没有让hipcc做device code的fat binary链接，最终可执行文件里可能就缺少GPU代码段。解决办法就是用hipcc做最终链接，并显式加上--hip-link。
6.3 分离编译后kernel运行结果不对是为什么
这个问题比较隐蔽。多数情况下是因为分离编译后，host代码和device代码的编译参数不一致，导致一些宏定义或类型元数据对不上。比如，头文件里有一个#ifdef控制的配置项，在编译color_convert.cpp时开启，编译main.cpp时却关闭，最终代码里两个模块对同一块内存的解释就不一致，kernel执行出来的结果自然不对。
处理思路是：把公共的编译宏（-D参数）统一提取到CMake的全局配置或统一的config文件里，让所有编译单元用同一组宏定义。同时，尽量控制头文件不要被不同的编译单元以不同的宏状态包含。这种问题排查起来很费时间，但掌握了以后会形成一个习惯：编译配置的一致性永远优先考虑。
另一个容易被忽视的点是hipDeviceSynchronize()的位置。分离编译后，kernel的启动可能被编译成异步的，如果不做同步就读取dst数据，读到的可能是没更新的脏数据。这本身和分离编译没有直接关系，但分离编译后代码更分散，更容易忘记加同步。
6.4 为什么分离编译后二进制变得更大了
有次我把一个算法库改成分离编译后，发现最终可执行文件的体积比之前全量编译增大了。一开始有点疑惑，后来查了fat binary的结构才明白：分离编译后，每个编译单元里只要有__global__函数，都会生成一份完整的device code object；而全量编译时，同一个编译单元内多个kernel共享了一些公共的device函数和类型信息。当你把多个文件拆开编译，这些公共部分就会在每个目标文件里各存一份，最终链接的时候，除非做重复代码合并，否则二进制体积确实可能变大。
解决办法是开启LTO（链接时优化）。HIP支持--hip-link配合LTO参数，让链接阶段的LLVM可以跨编译单元去重、合并和优化。但这会增加链接时间，增大内存消耗，适合在release构建时开启。我个人的经验是：调试阶段用分离编译不带LTO，发布前开LTO做一次全量重编，兼顾构建速度和最终产物体积。
6.5 如何确认GPU code object是否已经正确嵌入
有一次遇到kernel能编译但跑不起来的诡异情况，我手动检查了可执行文件里的code object内容。用llvm-objdump可以轻松查看：
如果输出有对应的GPU汇编，说明GPU代码确实存在于二进制中。如果提示找不到该符号，那你就要回溯编译流程，看看是不是device pass被跳过了。另一种方式是用readelf或objdump列出二进制段：
正常情况下你应该能看到类似.amdgcn或.text中包含AMDGCN code object的段。这个方法对于排查“编译通过但设备端啥也没跑”的问题非常高效。
7. 结合Makefile的轻量级分离编译实践
如果你项目还没引入CMake，或者只是临时做一次构建优化，直接用Makefile也能把分离编译跑起来。写一个简化版本的Makefile：
这个Makefile的核心点在于：color_convert.cpp和convolution.cpp用hipcc编译，生成带GPU code的目标文件；main.cpp纯host逻辑，用普通clang++编译。通过-MMD -MF生成依赖文件，Makefile通过-include读取它们，实现头文件驱动的增量构建。
这种轻量方案特别适合快速原型验证和内部工具链。真要上规模，还是建议用CMake，它对HIP语言有原生支持，依赖管理、并行编译、测试集成都更成熟。
8. 深入一点：__device__代码的LTO和跨模块优化
分离编译到这一步，基本已经解决了构建效率问题。但如果你追求极致性能，还想再进一步——把不同编译单元里的__device__函数做跨模块优化。
默认情况下，分离编译后每个模块的device代码是独立优化的。假设convolution.cpp里调用了一个在color_convert.cpp中定义的设备端辅助函数__device__ inline float clamp01(float v)，如果不做链接时优化，这个调用可能无法被完全内联展开，性能会受一点影响。
HIP的完整链路其实支持把device代码提取成bitcode再做LTO。思路是把所有源文件先编译成.bc格式，然后统一用llvm-link合并，再做一轮优化和代码生成。实际操作可以用hipcc的-c -emit-llvm来生成bitcode：
合并：
最后生成可执行文件或对象文件时再交给clang加上amdgcn的target参数。这个流程能做到跨文件的内联优化，但也会显著增加构建时间，适合在release版本上开。日常开发调试阶段，建议还是用基础的目标文件分离编译，迭代效率优先。
9. 从全量编译到分离编译：一次真实的构建优化记录
最后分享一个真实的数据。我之前维护过一个中型计算机视觉库，大概40个源文件，全量编译在16核的机器上约需12分钟。改成分离编译+增量构建之后，一次全量首编大概9分多钟（有少量中间文件的检查开销），但日常改一个函数引发的增量重建基本控制在1分钟以内，头文件大改时也只要3到5分钟。
我把要点总结成几条特别实在的经验：
- 源文件划分要精细，别把几十个函数塞进一个大文件里，否则改了哪都是全量重编；
- 头文件要稳定，能吃tolong的地方别自己写结构体，头文件变化是增量构建最大的敌人；
- GPU架构参数要在整个工程里统一，别让两个模块编成不同的架构，不然后期排查能让你怀疑人生；
- 分离编译并不复杂，本质就是CPU端早就用了几十年的“分而治之”思路，只是GPU编译链路更繁杂，需要多留意几个环节。
这套流程跑通以后，你会发现不止构建时间下来了，日常改代码的心态也稳了。不用每次编译都担心全量重来，改一行、编一次、立刻验证，这大概就是开发效率最直接的提升。后面如果你的项目持续变大，还可以考虑把device代码进一步抽成公共库、接入分发编译、利用Ninja并行编译，方向都是通的。
277 220 625 649 563 870 193 418 685 781 3681 884 1807 42 1305 818 1479 269 81 634 7 3