踩坑 PyTorch 三方库 CUDA → 国产加速卡适配？FastPT 一键搞定，不用肝 HIP 源码

接手过一个基于 PyTorch 开发的三方组件项目，里面塞了一堆 .cu 文件和 nvcc 编译指令。老板说"一周内把这个项目跑到国产加速卡上"。当时第一反应是：难道要一个一个算子手写 HIP？

后来发现 FastPT 这个工具，两条路可以走——不转码直接编译，或者转码把 CUDA 源码自动翻成 HIP。这篇文章就聊聊这两种路子的实际踩坑过程，以及那些不翻源码根本不知道的坑。

---

1. 先搞懂背景：CUDA 和国产加速卡的软件栈长什么样

聊工具之前，先看一眼软件栈。国产加速卡（DCU 架构，基于 GPGPU）和 NVIDIA GPU 软件栈的结构很像，差异主要在中间层：

一句话总结：国产加速卡的 ROCm/DTK 栈提供了 HIP 编程接口，相当于 CUDA 的"方言版"。FastPT 就是帮你处理这个方言转换的工具——要么直接用 GPUFusion 兼容层（不转码），要么把 CUDA 源码自动转成 HIP（转码）。

---

2. FastPT 是什么？一张图看懂

FastPT 是一个基于 Python 的应用编译工具，专门解决"基于 PyTorch + CUDA 代码的三方组件怎么在国产加速卡上编译运行"的问题。

GitHub 上的开源项目（如 torchvision、torchaudio、Megatron-LM、metaseq 等），大部分基于 CUDA 实现自定义算子和扩展。当你发现项目里没有 HIP/ROCm 分支，直接 clone 下来在国产加速卡环境里 python setup.py install 大概率会炸——FastPT 就是来解决这个的。

---

3. 选哪条路？不转码 vs 转码决策树

先看一个简单决策：

核心原则：能用不转码就不用转码。转码虽然自动化程度高，但遇到 warp 原语、Blas 类型转换等坑，改起来比不转码折腾得多。后面会细说。

---

4. 不转码编译实战：metaseq 亲测跑通

4.1 不转码编译流程

4.2 实操步骤

以 metaseq（Facebook Research 的大规模序列数据处理库）为例：

Step 1：对齐版本

组件	要求
DTK	25.04.x 系列
PyTorch	2.x 系列（最低 2.4.1）
FastPT	2.x 系列（与 DTK/PyTorch 版本对应）

FastPT 1.x 只支持转码，2.x 才开始支持不转码。注意版本对应关系，混搭会翻车。

Step 2：初始化编译环境

which fastpt
source /path/to/fastpt -C

输出大概长这样：

current_dtk_version: 25.04.x
Success: The current torch version matches the required version.
USE FASTPT CUDA is set, and CUDA environment is loaded

就这一行命令，FastPT 在背后干了三件事：

1. 初始化 GPUFusion 环境（source /opt/dtk/cuda/env.sh）
2. 给 torch 补充编译需要的头文件和动态库（比如 torch_cuda 的 .so）
3. 设置编译环境变量

Step 3：正常编译

pip install cython iopath
python setup.py bdist_wheel && pip install dist/metaseq*.whl --force --no-deps

Step 4：验证

python -c "import metaseq; print(metaseq.__version__)"

如果报 libc10_cuda.so: cannot open shared object file，执行一次 source /path/to/fastpt -E 就好。

4.3 踩坑记录

1. -C 和 -E 别混用。-C 初始化了 GPUFusion 环境，老版本 DTK 不支持取消 GPUFusion 设置，要测别的东西就新开终端。
2. C++ 标准要求 C++17 或以上。
3. __CUDA_ARCH__ 宏问题：老版本 DTK 下有些 __device__ 接口用 __CUDA_ARCH__ 控制条件编译，导致 host 端代码看不到函数定义。用 __CUDACC__ 替换即可，新版本已修复。
4. 别用 nv 源码编译的库：优先去 DAS 社区找适配好的版本，否则会遇到 CUDAHooksMocker 重复注册问题。
5. 静态库链接不支持：老版本 DTK 不支持 -lcudart_static，改成 -lcudart 动态链接。

---

5. 转码编译实战：metaseq 踩坑全记录

5.1 转码编译流程

5.2 实操步骤

Step 1：-T 转码编译

source /path/to/fastpt -T
python setup.py install

执行完后，源码目录下会自动生成 .hip 后缀的转码文件。比如：

metaseq/modules/apex/
  fused_adam_cuda_kernel.cu    ← 原始 CUDA 文件
  fused_adam_hip_kernel.hip    ← 自动转码生成的 HIP 文件
  fused_dense_cuda.cu
  fused_dense_cuda.hip

但是，直接编译大概率会炸。下面是实际踩过的几个坑。

---

5.3 坑一：__shfl_xor_sync 洗牌函数

现象：编译时报 error: use of undeclared identifier '__shfl_xor_sync'。

原因：国产加速卡下 __shfl_xor_sync 接口需要加编译选项 -DHIP_ENABLE_WARP_SYNC_BUILTINS 且 mask 设为 uint64 类型，或者直接换成不带 sync 后缀的 __shfl_xor。

修法：找到报错对应的原始 CUDA 源文件（不是 .hip 转码文件，因为重新转码会覆盖你的修改），加条件编译：

template <typename T>
__device__ __forceinline__ T WARP_SHFL_XOR_NATIVE(T value, int laneMask, 
    int width = warpSize, unsigned int mask = 0xffffffff) {
#if CUDA_VERSION >= 9000 && !defined(__HIP_PLATFORM_HCC__)
    return __shfl_xor_sync(mask, value, laneMask, width);
#else
    return __shfl_xor(value, laneMask, width);  // HIP 下走这条路
#endif
}

同理，__shfl_sync → __shfl，__shfl_up_sync → __shfl_up，__shfl_down_sync → __shfl_down，syncwarp() 直接删掉就行。

---

5.4 坑二：Blas 数学库类型对不上

现象：error: no matching function for call to 'hipblasGemmEx'，提示 no known conversion from 'hipDataType' to 'hipblasDataType_t'。

原因：转码工具把 CUDA 的 CUDA_R_xxx 自动转成了 HIP_R_xxx（即 hipDataType 类型），但 hipblasGemmEx 是 hipBLAS 接口，它要的参数类型是 hipblasDataType_t（即 HIPBLAS_R_xxx）。

搞混了三个类型：

类型	用于	枚举值示例
hipDataType	hipBLASLt 接口	HIP_R_32F
hipDataType_t	hipBLASLt 接口	HIP_DATATYPE_R_32F
hipblasDataType_t	hipBLAS 接口	HIPBLAS_R_32F

修法：在原始 CUDA 源文件中，把 cublasGemmEx 调用里的 CUDA_R_xxx 手动改成 HIPBLAS_R_xxx：

// 改前（CUDA 源码，自动转码后会变成 HIP_R_64F，类型不对）
return cublasGemmEx(handle, transa, transb, m, n, k, 
    alpha, A, CUDA_R_64F, lda, B, CUDA_R_64F, ldb,
    beta, C, CUDA_R_64F, ldc, CUDA_R_64F, CUBLAS_GEMM_DEFAULT);

// 改后（直接把类型写成 HIPBLAS_R_64F，转码后就是对的）
return cublasGemmEx(handle, transa, transb, m, n, k,
    alpha, A, HIPBLAS_R_64F, lda, B, HIPBLAS_R_64F, ldb,
    beta, C, HIPBLAS_R_64F, ldc, HIPBLAS_R_64F, CUBLAS_GEMM_DEFAULT);

---

5.5 坑三：接口重复定义

现象：编译报 error: use of overloaded operator '*' is ambiguous，比如 float2 和 float 的乘法运算符在两个头文件里都有定义。

修法：HIP 下如果已经有了相关实现，用条件编译把自定义实现包起来：

#ifndef __HIP_PLATFORM_HCC__
// 只在非 HIP 环境下生效的自定义实现
__device__ inline float2 operator-(const float2& a, const float2& b) {
    return make_float2(a.x - b.x, a.y - b.y);
}
__device__ inline float2 operator*(const float a, const float2& b) {
    return make_float2(a * b.x, a * b.y);
}
#endif

---

5.6 坑四：__device__ 属性丢失

有些模板函数在 CUDA 源码里没加 __device__ 修饰，在 CUDA 环境下能编译过，但转码到 HIP 后 host 端编译会报错。比如：

// 改前：rsqrt 模板函数缺少 __device__
template<typename U> U rsqrt(U v) { return U(1) / sqrt(v); }

// 改后：加上 __device__ 修饰
template<typename U> __device__ U rsqrt(U v) { return U(1) / sqrt(v); }
template<> __device__ float rsqrt(float v) { return rsqrtf(v); }
template<> __device__ double rsqrt(double v) { return rsqrt(v); }

---

5.7 转码接口速查

FastPT 提供了 Python 层面的 hipify 接口：

from torch.utils.hipify import hipify_python

hipify_python.hipify(
    project_directory=extensions_dir,      # 源码路径
    output_directory=extensions_dir,       # 输出路径
    includes='src/*',                      # 需要转换的目录
    is_pytorch_extension=True,             # PyTorch 扩展场景
    add_dtk_macros=True,                   # 是否添加 DTK 宏
    keep_file_path=False,                  # 是否保留原路径结构
)

对于 CMake 构建的项目，keep_file_path=True 可以保留源码目录结构，生成一个 xxx_dtk 的转码副本文件夹。

FastPT 还支持自定义转码映射：在同级目录下创建 custom_hipify_mappings.json，自定义 CUDA→HIP 的符号对应关系，优先级高于内置映射。

---

6. 不转码 vs 转码：一张表说清楚

维度	不转码（推荐）	转码
原理	GPUFusion 兼容层直接编译 CUDA 源码	CUDA 源码自动转 HIP → hipcc 编译
适用场景	绝大多数 PyTorch 三方组件	GPUFusion 不支持但 HIP 支持的功能（如 cutlass → hytlass）
上手难度	低，一行 -C 搞定	高，需要逐坑手动修源码
稳定性	高	中（warp 原语、Blas 类型等需要人工介入）
代码侵入	几乎零	需要修改源文件（加条件编译、替换类型等）
CMake 项目	可能需调整编译选项	需要手动转义 CMake 脚本
部署方式	-E 初始化运行环境	编译后直接运行，无需额外环境

一句话：能用不转码就别用转码。

---

7. 踩坑后总结的避坑指南

1. 版本对齐是第一要务：DTK、PyTorch、FastPT 三个版本必须对得上，去 DAS 社区的 FastPT 仓查版本对应表。
2. -C 和 -E 各司其职：-C 管编译，-E 管运行。编译好的 whl 包在新环境部署时只需 -E，别手贱再跑一次 -C。
3. 第一个 error 是关键：编译报错一堆时，优先解决第一个 error。后续的错往往是第一个引起的连锁反应。
4. 改源码不改变转码后的 .hip 文件：转码是覆盖式的，你改 .hip 文件下次转码就没了。定位到原始 .cu/.h 文件改，加条件编译兼容两边。
5. DAS 社区有现成的：FastPT 仓（developer.sourcefind.cn/codes/OpenDAS/fastpt）里汇总了大量已适配的应用，带 -fastpt 后缀的分支就是现成方案，别从零造轮子。
6. GPUFusion 不支持的功能怎么办：临时屏蔽 + 去 DAS 找替代库。实在绕不开再考虑转码。
7. 环境隔离：老版本 DTK 不支持 GPUFusion 环境取消，测试不同方案时开新终端。

---

8. 总结

FastPT 把 PyTorch 三方组件从 CUDA 到国产加速卡的适配门槛降了一大截。大部分场景走不转码编译，一行 source fastpt -C + 正常编译就搞定。只有在 GPUFusion 确实不支持、而 HIP 下有替代方案时，才上转码——但要做好和 __shfl_xor_sync、Blas 类型转换、接口重复定义这些坑死磕的心理准备。

回头看看，最省事的策略是：先查 DAS 上有没有现成的适配分支，没有再上 FastPT 不转码，最后再考虑转码。