# 1 FastPT简介 FastPT是基于python的应用编译工具，借助FastPT，开发人员可以在HCU上开发、部署基于pytorch的内含CUDA代码的应用。可以实现CUDA源码不转码直接编译，或源码转换到HIP格式代码后，通过hipcc适配编译。推荐优先使用不转码编译方式。转码适配方式功能受限，可能需要手动处理较多的内容。 FastPT版本与torch版本对应关系为： | | FastPT版本 | torch版本 | DTK版本 | | - | ----------------- | --------- | ------- | | 1 | 2.1.0+das.dtk2504 | v2.5.1 | dtk2504 | | 2 | 2.0.1+das.dtk2504 | v2.4.1 | dtk2504 | 注： 1. FastPT工具目前尚不支持cutlass类似三方库和内联汇编指令； 2. 编译有依赖库需求且依赖库涉及到GPU加速实现的，可以在DAS上以及GPUFusion（一般为$ROCM_PATH/cuda/)环境下查询是否有已适配或支持的库； 3. 此工具适合生态组件应用依赖torch的且有GPU加速代码实现的场景下使用，如通过CUDAExtension构建编译或编译依赖libtorch等。 # 2 安装 工具安装使用pip方式，安装前请确保环境中已安装了torch，并从光源社区-DAS中下载此工具的安装包。注意与python、torch版本匹配。Torch需要使用HCU下支持的版本。 ``` cd path/to/whl pip install fastpt*.whl --no-deps ``` 安装完成之后，可通过以下指令验证是否安装成功,指令执行后会显示当前fastpt的版本号。 ``` python -c "import fastpt;print(fastpt.__version__)" ``` # 3 使用 推荐使用不转码编译方式，可参考下面的表格以及3.1章节的内容。 工具安装后，**首先通过指令 which fastpt 来获取 fastpt 的安装路径（下面以 /usr/local/bin/fastpt 路径进行说明）**。在构建编译或使用时，通过source /usr/local/bin/fastpt -X 进行环境设置。X为模式设置参数，具体参数说明如下： | 使用场景 | 指令 | 示例 | 说明 | | ---------- | -------- | ------------------------------- | ------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------ | | 不转码编译 | -C 或 -c | source /usr/local/bin/fastpt -C | 用于工程不转码编译处理的环境设置,由于在编译模式下需要设置部分环境变量，所以在打开新的终端进行编译处理时，需要执行此命令。 | | 不转码编译 | -E 或 -e | source /usr/local/bin/fastpt -E | 不转码编译的程序执行时环境设置。用于工程在不转码编译后，进行使用时的环境设置。工程迁移到新环境后，安装FastPT后，执行此命令即可。此命令不需要重复执行，只需保证当前系统下执行过即可。 | | 转码编译 | -T 或 -t | source /usr/local/bin/fastpt -T | 转码编译处理。用于通过转码方式，将CUDA代码转换到HIP代码后的编译实现。只用于组件或程序编译处理，组件执行时不需要额外配置环境。 | | 帮助 | -H 或 -h | source /usr/local/bin/fastpt -H | 具体指令说明，查询使用方法。 | ## 3.1 不转码编译 ### 3.1.1 使用方法 编译模式：Fastpt-2.0之后支持不转码编译实现，即直接使用CUDA源码编译。安装FastPT工具的whl包，执行： ``` source /usr/local/bin/fastpt -C ``` 初始化FastPT编译环境，然后按照组件或应用的官方指导编译方法处理即可。 执行模式：编译好组件通过whl包在有FastPT的新环境下使用时，需要确保环境下已执行过执行模式的初始化操作，否则可能会报错找不到CUDA相关的动态库。执行命令： ``` source /usr/local/bin/fastpt -E ``` 即可。 ### 3.1.2 不转码编译示例 下面以pytorch_scatter组件为例，通过fastpt不转码编译方式来编译该组件。 首先，启动一个容器，对当前容器内的环境进行下列检查： （1）是否安装DTK且DTK的版本是否正确； （2）是否安装pytorch且pytorch的版本是否正确； （3）是否安装FastPT且FastPT的版本是否正确。 检查完毕后，下载 pytorch_scatter 组件源码，并执行下面的命令进入编译模式： ``` source /usr/local/bin/fastpt -C ``` * 编译模式: 进入编译模式后，会有如下提示： ``` current_dtk_version: 2504 WARNING: The current dtk version is consistent with the required dtk version. Success: The current torch version 2.4 matches the required version 2.4. Torch version is correct: 2.4 torch version check success! USE_FASTPT_CUDA is set, and CUDA environment is loaded ``` 此时已初始化FastPT编译环境。 接着，查阅pytorch_scatter的readme.md文档，找到组件的官方指导编译方法。针对该组件，运行下面的命令执行编译并生成whl包： ``` python setup.py bdist_wheel ``` 编译完成后，生成的whl包会在dist目录下。可以执行下面的命令安装pytorch_scatter并自动安装其依赖包： ``` pip install torch_scatter-2.1.0-cp310-cp310-linux_x86_64.whl ``` * 执行模式: whl包安装完成后，为测试编译的组件是否可用，执行下面的命令进入执行模式： ``` source /usr/local/bin/fastpt -E ``` 此时终端会有如下提示： ``` current_dtk_version: 2504 WARNING: The current dtk version is consistent with the required dtk version. No ROCm runtime is found, using ROCM_HOME='/opt/dtk' Success: The current torch version 2.4 matches the required version 2.4. Torch version is correct: 2.4 torch version check success! torch mocker so already exist current_dtk_version: 2504 fastpt.cuda has been initialized ``` 大部分组件内通常会存在测试用例用于验证，可以通过执行下面的命令进行组件测试： ``` pytest -vs ./tests ``` 测试完成后，如果用例全部成功测试通过，则该组件验证可用。如果出现失败用例，可以分析其失败原因，从而验证该组件是否可用。 ### 3.1.3 不转码编译注意事项 （1）此工具适用于依赖torch的生态组件使用，应用中内含CUDA C/C++代码的工程在HCU环境下的开发、移植。Fastpt版本应与torch版本对应，当前torch版本要求最低torch-2.4.1； （2）不支持依赖cutlass、内嵌汇编代码的编译； （3）使用__CUDA_ARCH__实现条件编译，当使用此宏控制__device__接口是否可见时，由于该宏作用于device代码的编译，会导致接口声明或定义不可见，建议使用__CUDACC__来实现条件编译，举例说明如下，在 host 端代码编译的时候，编译器会严格执行 host 端有关__global__等函数的语法检查，导致被__CUDA_ARCH__包裹的 reduceAdd 函数不可见： ```C++ #ifdef __CUDA_ARCH__ // 该处 __CUDA_ARCH__ 在编译 host 端代码时导致 reduceAdd 函数不可见 -> 报错信息：error: use of undeclared identifier 'reduceAdd' __device__ __forceinline__ static void reduceAdd(float *address, float val) { #if (__CUDA_ARCH__ < 200) // 该处 __CUDA_ARCH__ 在编译 host 端代码时无影响 #ifdef _MSC_VER #pragma message( \ "compute capability lower than 2.x. fall back to use CAS version of atomicAdd for float32") #else #warning \ "compute capability lower than 2.x. fall back to use CAS version of atomicAdd for float32" #endif int *address_as_i = reinterpret_cast(address); int old = *address_as_i, assumed; do { assumed = old; old = atomicCAS(address_as_i, assumed, __float_as_int(val + __int_as_float(assumed))); } while (assumed != old); #else atomicAdd(address, val); #endif } #endif // __CUDA_ARCH__ ``` （4） 编译模式(-C)与执行模式(-E)下，torch.version.cuda与torch.version.hip会分别被设置，少部分应用在执行时会依赖这两个变量，需要根据具体情况在应用端调整上述两个变量。 （5） 编译时遇到 `fatal error: #include_next ` 的报错，可以降低CMake版本为3.19处理。 ## 3.2 转码编译 FastPT提供了HCU下，转码到HIP格式，通过hipcc进行编译的方法，实现基于torch的应用中CUDA代码移植到HCU平台，工具接口包括CUDAExtension、CppExtension、hipify转码接口。编译时，转码一般是自动实现的。另外提供了自定义接口映射用来补充代码映射关系；提供了保持源码文件夹下文件相对路径的转码方法。 ### 3.2.1 使用方法 此方法适用于通过setup.py使用CUDAExtension、CppExtension进行组件构建编译的场景。使用时，执行 ``` source /usr/local/bin/fastpt -T ``` ### 3.2.2 自定义接口映射 工具中可能存在未涉及到或用户需要的一些转换匹配，可以通过json文件的方式给到工具，在不需要额外修改代码的情况下，实现自定义代码匹配转换。可以通过以下方法补充代码映射。用户需要在CUDAExtension或hipify_python接口调用代码同级目录下，通常为setup.py文件所在目录，创建 custom_hipify_mappings.json文件。json文件内容格式如下: ``` { "custom_map" : { "src mapping 1" : "dst mapping 1", "src mapping 2" : "dst mapping 2", ... } } ``` 此示例中，在将CUDA代码转到HIP代码时，会将源码中的”src mapping 1”替换为”dst mapping 1”，将"src mapping 2" 替换成 "dst mapping 2"。 自定义映射转换优先级高于内置的转换。 ### 3.2.3 转码接口 工具提供了hipify转码接口，使用可参考torch中的同名接口。 此接口还提供了实现保留源代码路径的转码的处理方法，转码会新建一个xxx_dtk的代码文件夹，内部文件路径结构与源文件夹下一致，".cu"代码文件扩展名会转成".hip"，如xxx.cu转码后为xxx.hip。此方法一般在CMake或Make等需要保持原代码路径的工具编译代码时使用，需要注意CMakeLists.txt中的代码需要手动修改，暂时不支持CMake语法的转码，如".cu"后缀需要改为".hip"、编译器nvcc需要改为hipcc等，此适配场景下建议使用3.1章节中的不转码编译处理。 使用时，需要在要转码的文件夹的同级目录中实现以下脚本fastptcode.py： ```python import os this_dir = os.path.dirname(os.path.abspath(__file__)) from fastpt import hipify res = hipify( project_directory=os.path.join(this_dir,'codepath'), includes = '*', show_detailed=True, is_pytorch_extension=True, add_dtk_macros=False, # False: 不引入ATen/dtk_macros.h 头文件，默认为True keep_file_path=True # True:保持源代码文件夹内文件相对路径;默认False # ignores=["run_tests.sh.in"] # 屏蔽掉不希望处理的代码 ) ``` 执行: ``` python fastptcode.py ``` 会在codepath的同级目录下生成转码后的文件夹 codepath_dtk，此路径下的文件结构与codepath的结构相同，”.cu“代码文件后缀变为”.hip“。文件结构如下： ``` | |----codepath |----codepath_dtk |----fastptcode.py ``` ### 3.2.4 注意事项 （1）此工具适用于依赖torch的生态组件或应用，内含CUDA C/C++代码的工程在HCU环境下的开发、移植，注意FastPT版本与torch版本对应; （2）暂不支持CMake、make等代码语义处理。代码转换可通过上面3.2.3中的示例，通过执行python的转码脚本代码，将CUDA代码转换成HIP代码，CMake文件需要用户自行处理。为了便于适配建议通过CMake编译的组件使用3.1章节的不转码的方式，通过CUDA源码编译的方式适配; （3）工程中存在三方依赖库时，三方库可能存在不被处理的情况，此时需要对三方依赖库进行单独处理; （4）适配组件的setup.py 中可能会有CUDA环境检查来决定是否执行CUDA相关代码的编译，例如CUDA_PATH或torch.version.cuda的检查，可按情况进行处理; （5）当不希望引入ATen/dtk_macros.h这个头文件时，可以通过以方式屏蔽此头文件的引入：export FASTPT_DTK_MACROS=1。 # 4 已支持组件列表 在 OpenDAS 仓库中，分支名带有 fastpt 后缀的组件支持 FastPT 不转码编译，具体如下： | 组件名称 | 版本及分支 | DAS仓库 | 已验证的FastPT版本 | | :-----------------: | :--------------------: | ------------------------------------------------------------------ | --------------- | | audio | v2.4.1-fastpt; v2.5.1-fastpt | https://developer.sourcefind.cn/codes/OpenDAS/torchaudio |2.0.1；2.1.0 | | vision | v0.19.1-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/vision.git |2.0.1；2.1.0 | | mmcv | v2.1.0-fastpt; v2.2.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/mmcv.git | 2.0.1；2.1.0 | | pytorch3d | V0.7.8-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/pytorch3d.git |2.0.1；2.1.0 | | maskrcnn | v0.1-fastpt | https://developer.sourcefind.cn/codes/OpenDAS/maskrcnn | 2.0.1；2.1.0 | | mmdetection3d | v1.4.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/mmdetection3d.git | 2.0.1；2.1.0 | | opencv-python | 4.8.0 | http://developer.sourcefind.cn/codes/OpenDAS/opencv-python.git | | | pytorch_sparse | 0.6.16-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/torch-sparce.git |2.0.1；2.1.0 | | pytorch_scatter | 2.1.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/torch-scatter.git |2.0.1；2.1.0 | | pytorch_cluster | 1.6.3-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/torch-scatter.git |2.0.1；2.1.0 | | pytorch_spline_conv | 1.2.2-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/torch-spline-conv.git |2.0.1；2.1.0 | | torchnai | v2.2.4-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/torchani.git | 2.0.1；2.1.0 | | apex | 24.04.1-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/apex.git |2.0.1；2.1.0 | | fastmoe | v1.1.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/fastmoe.git | 2.0.1；2.1.0 | | lietorch | v0.3-fastpt | https://developer.sourcefind.cn/codes/OpenDAS/lietorch | 2.0.1；2.1.0 | | uni-core | v0.0.1-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/Uni-Core.git |2.0.1；2.1.0 | | OpenPCDet | v0.6.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/openpcdet.git |2.0.1；2.1.0 | | SparseConvNet | v0.2-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/sparseconvnet.git |2.0.1；2.1.0 | | detectron2 | v0.6-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/detectron2.git |2.0.1；2.1.0 | | fairscale | v0.4.9-fastpt、v0.4.3-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/fairscale.git | 2.0.1；2.1.0 | | fairseq | v0.9.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/fairseq.git |2.0.1；2.1.0 | | metaseq | main-f7ffa5fd-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/metaseq.git | 2.0.1；2.1.0 | | pydensecrf | master-2723c7fa-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/pydensecrf.git | 2.0.1；2.1.0 | | d2go | mian-f4ac1567-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/d2go.git |2.0.1；2.1.0 | | dlib | v19.24-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/dlib.git | 2.0.1；2.1.0 | | causal_conv1d | v1.5.0-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/causal-conv1d.git |2.0.1；2.1.0 | | cubvh | main-ee89d5fa-fastpt | http://developer.sourcefind.cn/codes/OpenDAS/cubvh.git |2.0.1；2.1.0 | # 5 附录 fastpt whl包链接地址：https://download.sourcefind.cn:65024/4/main/fastpt