一、Transition Tool: HIPIFY
-
将 hip/bin 路径添加到 PATH:这一步确保
hipify-perl
和hipcc
(HIP 编译器)命令在您的命令行环境中可用。[MYHIP]
应该替换为您 HIP 安装目录的实际路径。 -
定义环境变量:设置
HIP_PATH
环境变量,让系统知道在哪里可以找到 HIP 相关的文件和库。 -
构建可执行文件:
- 切换到示例代码所在的目录。
- 运行
make
,这会触发hipify-perl
脚本将 CUDA 代码(square.cu
)转换为 HIP 代码(square.cpp
)。 - 使用
hipcc
将 HIP 代码编译成可执行文件。创建了两个版本:一个使用动态链接(square.out
),另一个使用静态链接(square.out.static
)。
-
执行文件:运行可执行文件将执行 HIP 代码,输出显示了它运行的设备信息、内存分配、内核启动和计算结果。
二、HIP Porting Process
1. Porting a New CUDA Project
General Tips
以下是一些将 CUDA 代码移植到 HIP 的通用提示:
-
在 CUDA 机器上开始移植通常最简单:您可以逐步将代码片段移植到 HIP,同时保留其余部分在 CUDA 中。(请记住,在 CUDA 机器上,HIP 只是 CUDA 上的一层薄层,因此两种代码类型可以在 nvcc 平台上互操作。)此外,HIP 移植版本可以与原始 CUDA 代码进行功能和性能比较。
-
一旦 CUDA 代码移植到 HIP 并在 CUDA 机器上运行:在 AMD 机器上使用 HIP 编译器编译 HIP 代码。
-
HIP 移植可以替代 CUDA 版本:HIP 可以提供与原生 CUDA 实现相同的性能,并且具有可移植性到 Nvidia 和 AMD 架构的优势,以及未来 C++ 标准支持的路径。您可以通过条件编译或将其添加到开源 HIP 基础设施中来处理平台特定的特性。
-
使用 bin/hipconvertinplace-perl.sh 工具 hipify 所有 CUDA 源代码目录中的代码文件:这个脚本可以帮助您自动化转换过程,将整个目录中的 CUDA 代码批量转换为 HIP 代码。
这些提示可以帮助您更高效地将 CUDA 代码迁移到 HIP,以便在不同的 GPU 架构上实现代码的可移植性和性能优化。
Scanning existing CUDA code to scope the porting effort
hipexamine-perl.sh
是一个用于分析 CUDA 代码的工具,它可以扫描源代码目录,确定哪些文件包含 CUDA 代码以及这些代码中有多少可以自动转换为 HIP 代码。这个工具对于评估代码移植到 HIP 的可行性和工作量非常有用。
-
进入示例目录:
cd examples/rodinia_3.0/cuda/kmeans
这一步是切换到包含 CUDA 代码的目录。
-
运行 hipexamine-perl.sh 工具:
$HIP_DIR/bin/hipexamine-perl.sh
这一步是执行
hipexamine-perl.sh
工具,其中$HIP_DIR
应该替换为您的 HIP 安装目录。 -
工具输出:
- info: hipify 文件名:表示工具正在检查并尝试将指定文件中的 CUDA 代码转换为 HIP 代码。
- info: converted:显示转换的统计信息,包括设备(dev)、内存(mem)、内核(kern)、内置函数(builtin)、数学函数(math)、流(stream)、事件(event)、错误处理(err)、定义(def)、纹理(tex)和其他(other)的转换数量。警告(warn)和代码行数(LOC)也显示出来。
- info: TOTAL-converted:显示所有文件的总转换统计信息。
例如,输出中的一行:
info: converted 49 CUDA->HIP refs( dev:3 mem:32 kern:2 builtin:0 math:0 stream:0 event:0 err:0 def:0 tex:12 other:0 ) warn:0 LOC:311
表示在转换过程中,有 49 个 CUDA 引用被转换为 HIP,其中涉及 3 个设备引用、32 个内存引用、2 个内核引用、0 个内置函数引用等。没有警告,总共处理了 311 行代码。
-
内核统计:
- kernels (1 total):显示转换过程中发现的内核数量和名称,例如
kmeansPoint(1)
表示有一个名为kmeansPoint
的内核。
- kernels (1 total):显示转换过程中发现的内核数量和名称,例如
-
无 CUDA 代码的文件:
- 对于不含 CUDA 代码的文件(如
kmeans.h
),hipexamine-perl
仅打印一行摘要,列出源文件名。 - 例如:
info: hipify ./kmeans.h =====>
- 对于不含 CUDA 代码的文件(如
-
含 CUDA 代码的文件:
-
对于包含 CUDA 代码的文件(如
kmeans_cuda_kernel.cu
),工具会打印详细的摘要,包括:- 转换的 CUDA 调用数量:显示转换为 HIP 的 CUDA 调用总数。
- CUDA 功能使用情况的细分:详细列出了文件中使用的 CUDA 功能类别(如设备、内存、内核、内置函数、数学函数、流、事件等)。
- 警告:对于看似 CUDA API 但未被转换的代码,会给出警告。
- 代码行数(LOC):统计处理的代码行数。
-
例如:
info: hipify ./kmeans_cuda_kernel.cu =====> info: converted 40 CUDA->HIP refs( dev:0 mem:0 kern:0 builtin:37 math:0 stream:0 event:0 err:0 def:0 tex:3 other:0 ) warn:0 LOC:185
这表示在
kmeans_cuda_kernel.cu
文件中,有 40 个 CUDA 调用被转换为 HIP,其中使用了 37 个 CUDA 内置函数和 3 个纹理函数,没有警告,处理了 185 行代码。
-
-
过程结束时的总结:
-
hipexamine-perl
在分析完所有文件后,会提供一个总结,包括所有文件的统计数据。 -
总结的格式与单个文件报告类似,还包括所有已调用的内核列表。
-
例如:
info: TOTAL-converted 89 CUDA->HIP refs( dev:3 mem:32 kern:2 builtin:37 math:0 stream:0 event:0 err:0 def:0 tex:15 other:0 ) warn:0 LOC:3607 kernels (1 total) : kmeansPoint(1)
这表示在整个分析过程中,共转换了 89 个 CUDA 调用到 HIP,涉及 3 个设备、32 个内存、2 个内核、37 个内置函数等。没有警告,总共处理了 3607 行代码。此外,还列出了一个内核
kmeansPoint
。
-
使用场景
- 评估代码移植的可行性:通过分析报告,您可以了解代码中 CUDA 功能的分布和复杂性,从而评估将 CUDA 代码移植到 HIP 的难度和工作量。
- 优化移植策略:报告可以帮助您确定需要手动调整的代码部分,以及哪些部分可以自动转换。
- 监控移植进度:在移植过程中,您可以定期运行
hipexamine-perl
来监控进度和转换效果。
Converting a project in-place
> hipify-perl --inplace
hipify-perl --inplace
是一个用于将 CUDA 代码转换为 HIP 代码的命令行工具,它提供了一种就地(in-place)转换文件的方式。这个命令对于测试 hipify
工具集的改进非常有用,因为它允许您直接在原始文件上进行转换,同时保留原始代码的备份。以下是该命令的详细说明:
命令功能
- 备份原始代码:如果原始文件的
.prehip
备份版本不存在,该命令会首先将原始代码复制到一个带有.prehip
扩展名的新文件中。 - 执行转换:然后,它会将
.prehip
文件中的 CUDA 代码转换为 HIP 代码。 - 保存转换结果:如果
.prehip
文件已经存在,它会直接在该文件上进行 HIP 转换,并将结果保存回原始文件。
hipconvertinplace-perl.sh:这是一个脚本,用于在指定目录中的所有代码文件上执行就地转换。这个脚本会保留现有的目录结构和文件名,非常适合处理大型的 CUDA 代码库。
使用示例
假设您有一个名为 MY_SRC_DIR
的目录,其中包含需要转换的 CUDA 代码。您可以使用以下命令来执行就地转换:
hipconvertinplace-perl.sh MY_SRC_DIR
这个命令会遍历 MY_SRC_DIR
目录中的所有文件,并对每个文件执行 hipify-perl --inplace
操作。转换后,您可以检查转换结果,并根据需要添加额外的参数或进行进一步的手动调整。
注意事项
- 备份:在执行就地转换之前,确保您有原始代码的备份,以防转换过程中出现任何问题。
- 代码审查:转换后,仔细审查代码,确保所有 CUDA 特定功能都已正确转换为 HIP 代码,并且没有引入任何新的错误或性能问题。
- 测试:在实际部署之前,对转换后的代码进行充分的测试,以确保它在目标硬件上正常工作。
Library Equivalents
2. Distinguishing Compiler Modes
Identifying HIP Target Platform
在 HIP 编程环境中,正确识别和设置目标平台是确保代码能够正确编译和执行的关键步骤。HIP 支持两种主要的 GPU 平台:AMD 和 NVIDIA。以下是如何通过预定义宏来识别和区分这些平台:
目标平台宏定义
-
HIP_PLATFORM_AMD
- 当您的 HIP 项目目标是 AMD 平台时,编译器会定义
HIP_PLATFORM_AMD
。 - 这确保了在编译过程中包含的是适合 AMD GPU 的头文件和库。
- 当您的 HIP 项目目标是 AMD 平台时,编译器会定义
-
HIP_PLATFORM_NVDIA
- 当您的 HIP 项目目标是 NVIDIA 平台时,编译器会定义
HIP_PLATFORM_NVDIA
。 - 这确保了在编译过程中包含的是适合 NVIDIA GPU 的头文件和库。
- 当您的 HIP 项目目标是 NVIDIA 平台时,编译器会定义
编译和链接的影响
-
头文件的包含:不同的平台可能需要不同的头文件。例如,某些特定于平台的功能可能只在特定平台的头文件中定义。
-
库的链接:不同的平台可能需要链接不同的库。例如,AMD 和 NVIDIA 的 GPU 可能需要不同的数学库或驱动程序支持库。
#ifdef HIP_PLATFORM_AMD
// 包含 AMD 特定的头文件
#include <amd_specific_header.h>
#elif defined(HIP_PLATFORM_NVDIA)
// 包含 NVIDIA 特定的头文件
#include <nvidia_specific_header.h>
#else
// 通用代码或错误处理
#error "Unknown HIP platform"
#endif
Identifying the Compiler: HIP-Clang or NVIDIA
通过预定义的宏来判断编译器类型:
编译器宏定义
-
HIP_PLATFORM_AMD
- 当代码使用 HIP-Clang 编译器针对 AMD 平台编译时,会定义此宏。
- 这通常意味着代码正在使用 HIP-Clang,它是 AMD 提供的基于 LLVM/Clang 的编译器,用于 HIP 编程。
-
HIP_PLATFORM_NVIDIA
- 当代码使用 NVIDIA 的 nvcc 编译器针对 NVIDIA 平台编译时,会定义此宏。
- nvcc 是 NVIDIA 的 CUDA 编译器,它也可以用于编译 HIP 代码,因为 HIP 设计为与 CUDA 兼容。
CUDA 语言扩展宏
- CUDACC
- 当代码通过 nvcc 编译,并且启用了 CUDA 语言扩展时,会定义此宏。
- 这通常用于识别正在编译的是
.cu
文件,它允许使用 CUDA 特有的语言特性。
#ifdef __HIP_PLATFORM_AMD__
// Compiled with HIP-Clang
#endif
#ifdef __HIP_PLATFORM_NVIDIA__
// Compiled with nvcc
// Could be compiling with CUDA language extensions enabled (for example, a ".cu file)
// Could be in pass-through mode to an underlying host compile OR (for example, a .cpp file)
#ifdef __CUDACC__
// Compiled with nvcc (CUDA language extensions enabled)
注意事项:
- HIP-Clang:不定义
__CUDACC__
,因为它直接生成主机代码,而不需要传递给其他主机编译器。 - nvcc:当编译
.cu
文件时定义__CUDACC__
,允许使用 CUDA 特有的语言特性。
Identifying Current Compilation Pass: Host or Device
在 GPU 编程中,编译器如何处理代码对于性能优化和代码结构至关重要。NVIDIA 的 nvcc 编译器和 AMD 的 HIP-Clang 编译器在编译过程中有不同的行为模式。了解这些差异可以帮助开发者更有效地编写和优化 GPU 代码。
NVCC 的编译过程
-
两遍编译过程:
- 主机代码编译:首先编译所有不特定于设备的代码,即所有不在
__global__
函数或设备函数中的代码。 - 设备代码编译:然后编译所有在
__global__
函数或设备函数中的代码,针对特定的 GPU 架构。
- 主机代码编译:首先编译所有不特定于设备的代码,即所有不在
-
CUDA_ARCH:
- 这是一个在设备代码编译期间定义的宏,表示目标 GPU 的计算能力。例如,如果目标是计算能力为 75 的 GPU,则
__CUDA_ARCH__
会被定义为 75。
- 这是一个在设备代码编译期间定义的宏,表示目标 GPU 的计算能力。例如,如果目标是计算能力为 75 的 GPU,则
HIP-Clang 的编译过程
-
多遍编译过程:
- 主机代码编译:编译所有不特定于设备的代码。
- 设备代码编译:为每种目标设备架构编译设备代码。这可能包括为不同的 GPU 型号或计算能力编译不同的代码版本。
-
HIP_DEVICE_COMPILE:
- 这是一个在编译设备代码时定义的宏,无论是在
__global__
函数还是设备函数中。 - 与
__CUDA_ARCH__
不同,__HIP_DEVICE_COMPILE__
的值总是 1 或未定义,它不表示目标设备的功能能力。
- 这是一个在编译设备代码时定义的宏,无论是在
// #ifdef __CUDA_ARCH__
#if __HIP_DEVICE_COMPILE__
使用场景
- HIP_DEVICE_COMPILE:
- 用于确定代码是否正在作为设备代码编译,无论它是在 HIP-Clang 还是 nvcc 下。
- 可以替代对
__CUDA_ARCH__
的检查,因为后者仅在 nvcc 编译 CUDA 代码时有效。