README.md

# 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<int *>(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会分别被设置，少部分应用在执行时会依赖这两个变量，需要根据具体情况在应用端调整上述两个变量。

## 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 已支持组件列表

|      组件名称      |          版本          | DAS仓库                                                            | 已验证的FastPT版本  |
| :-----------------: | :--------------------: | ------------------------------------------------------------------ | --------------- |
|        audio        |     v2.4.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.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         | 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