HIP 支持¶
HIP (可移植异构计算接口) https://github.com/ROCm-Developer-Tools/HIP 是一个 C++ 运行时 API 和内核语言。它使开发人员能够从单个源代码创建可移植应用程序，以将计算卸载到不同的硬件平台。
AMD GPU 支持¶
Clang 通过 ROCm 平台 https://rocm.docs.amd.com/en/latest/# 在 AMD GPU 上提供 HIP 支持。ROCm 运行时构成了 HIP 主机 API 的基础，而 HIP 设备 API 则是通过 HIP 头文件和 ROCm 设备库实现的。Clang 驱动程序使用 HIPAMD 工具链通过 AMDGPU 后端将 HIP 设备代码编译为 AMDGPU ISA。然后将编译后的代码捆绑并嵌入到主机可执行文件中。
英特尔 GPU 支持¶
Clang 使用 CHIP-Star 项目 https://github.com/CHIP-SPV/chipStar 在英特尔 GPU 上提供部分 HIP 支持。CHIP-Star 在 oneAPI Level Zero 或 OpenCL 运行时之上实现了 HIP 运行时。Clang 驱动程序使用 HIPSPV 工具链将 HIP 设备代码编译成 LLVM IR，然后通过 SPIR-V 后端或树外 LLVM-SPIRV 翻译器将其翻译成 SPIR-V。然后将 SPIR-V 捆绑并嵌入到主机可执行文件中。
注意
虽然 Clang 没有直接为 NVIDIA GPU 和 CPU 提供 HIP 支持，但这些平台通过其他方式支持
NVIDIA GPU：HIP 支持通过 HIP 项目 https://github.com/ROCm-Developer-Tools/HIP 提供，该项目提供了一个仅限头文件的库，用于将 HIP 运行时 API 翻译成 CUDA 运行时 API。然后使用 NVIDIA 的 nvcc 编译代码。
CPU：HIP 支持可通过 HIP-CPU 运行时库 https://github.com/ROCm-Developer-Tools/HIP-CPU 获得。这个仅限头文件的库使 CPU 能够执行未修改的 HIP 代码。
示例用法¶
要编译 HIP 程序，请使用以下命令
clang++ -c --offload-arch=gfx906 -xhip sample.cpp -o sample.o
-xhip 选项表示源代码是 HIP 程序。如果文件具有 .hip 扩展名，Clang 将自动将其识别为 HIP 程序
clang++ -c --offload-arch=gfx906 sample.hip -o sample.o
要链接 HIP 程序，请使用此命令
clang++ --hip-link --offload-arch=gfx906 sample.o -o sample
在上面的命令中，--hip-link 标志指示 Clang 链接 HIP 运行时库。但是，如果您的程序中已经存在 HIP 输入文件，则不需要使用此标志。
为了方便起见，Clang 还支持在单个步骤中编译和链接
clang++ --offload-arch=gfx906 -xhip sample.cpp -o sample
在上面的命令中，gfx906 是正在为其编译代码的 GPU 架构。支持的 GPU 架构可以在 AMDGPU 处理器表 中找到。或者，您可以使用随 Clang 提供的 amdgpu-arch 工具来列出系统上的 GPU 架构
amdgpu-arch
您可以使用 --offload-arch=native 自动检测系统上的 GPU 架构
clang++ --offload-arch=native -xhip sample.cpp -o sample
依赖项路径设置¶
编译 HIP 程序取决于 HIP 运行时和设备库。HIP 运行时和设备库的路径可以通过编译器选项或环境变量指定。如果路径遵循 ROCm 安装目录结构，也可以通过 ROCm 路径设置它们。
HIP 路径优先级顺序¶
--hip-path编译器选项HIP_PATH环境变量 (谨慎使用)--rocm-path编译器选项ROCM_PATH环境变量 (谨慎使用)默认自动检测 (相对于 Clang 或在默认 ROCm 安装位置)
设备库路径优先级顺序¶
--hip-device-lib-path编译器选项HIP_DEVICE_LIB_PATH环境变量 (谨慎使用)--rocm-path编译器选项ROCM_PATH环境变量 (谨慎使用)默认自动检测 (相对于 Clang 或在默认 ROCm 安装位置)
编译器选项 |
环境变量 |
描述 |
默认值 |
|---|---|---|---|
|
|
指定 ROCm 安装路径。 |
自动检测 |
|
|
指定 HIP 运行时安装路径。 |
由 ROCm 目录结构确定 |
|
|
指定 HIP 设备库安装路径。 |
由 ROCm 目录结构确定 |
注意
我们建议使用编译器选项作为指定这些路径的主要方法。虽然环境变量 ROCM_PATH、HIP_PATH 和 HIP_DEVICE_LIB_PATH 受支持，但使用它们会导致隐式依赖，这可能会导致长期的问题。谨慎使用它们。
预定义宏¶
宏 |
描述 |
|---|---|
|
当 Clang 在可重定位设备代码 (RDC) 模式下编译代码时定义。RDC 由 |
|
在使用 HIP 语言支持进行编译时定义，表示代码针对 HIP 环境。 |
|
__HIP__ 的别名。 |
|
在 Clang 的主机和每个卸载 GPU 架构的单独编译过程中，在设备代码编译期间定义。 |
|
表示 HIP 中的单线程内存范围（值为 1）。 |
|
表示 HIP 中的波前内存范围（值为 2）。 |
|
表示 HIP 中的工作组内存范围（值为 3）。 |
|
表示 HIP 中的代理内存范围（值为 4）。 |
|
表示 HIP 中的系统范围内存范围（值为 5）。 |
|
当目标设备不支持 HIP 图像函数时，定义为 1。 |
|
__HIP_NO_IMAGE_SUPPORT__ 的别名。已弃用。 |
|
当 GPU 默认流设置为每线程模式时定义。 |
|
__HIP_API_PER_THREAD_DEFAULT_STREAM__ 的别名。已弃用。 |
注意，一些特定于体系结构的 AMDGPU 宏在从 HIP 主机编译时将具有默认值。其他 AMDGPU 宏 例如 __AMDGCN_WAVEFRONT_SIZE__ 将默认值为 64。
编译模式¶
每个 HIP 源文件都包含交织的设备和主机代码。根据编译器选项 -fno-gpu-rdc 和 -fgpu-rdc 选择的编译模式，这些代码部分的编译方式不同。
设备代码编译¶
``-fno-gpu-rdc`` 模式（默认）:
为每个卸载设备架构编译成一个自包含的、完全链接的卸载设备二进制文件。
翻译单元 (TU) 内的设备代码无法调用位于另一个 TU 中的函数。
``-fgpu-rdc`` 模式:
为每个 GPU 架构编译成一个位代码。
对于每个卸载设备架构，来自不同 TU 的位代码将链接在一起，以创建一个单个卸载设备二进制文件。
一个 TU 中的设备代码可以调用位于另一个 TU 中的函数。
主机代码编译¶
两种模式:
为每个 TU 编译成一个可重定位对象。
然后将这些可重定位对象链接在一起。
一个 TU 中的主机代码可以调用另一个 TU 中的主机函数并启动内核。
与 CUDA 的语法差异¶
Clang 的前端用于 CUDA 和 HIP 编程模型，共享相同的解析和语义分析机制。这包括关于设备和主机函数的重载解析。虽然在 Clang 和 NVCC 之间的方言差异 中存在关于 Clang 和 NVCC 之间 CUDA 语法差异的全面文档，但需要注意的是，这些差异也适用于 HIP 代码编译。
用于区分的预定义宏¶
为了便于区分 HIP 和 CUDA 代码，以及 HIP 中的设备和主机编译之间的区别，Clang 定义了特定的宏
__HIP__：仅在编译 HIP 代码时定义此宏。它可用于有条件地编译特定于 HIP 的代码，使开发人员能够编写可在 CUDA 和 HIP 两种环境下编译的可移植代码。__HIP_DEVICE_COMPILE__：仅在 HIP 设备编译期间定义，此宏允许有条件地编译特定于设备的代码。它提供了一种隔离设备和主机代码的机制，确保每个代码都可以针对各自的执行环境进行优化。
函数指针支持¶
在 Clang 与 HIP 中，函数指针的支持取决于使用模式。下表概述了不同用例和模式下的支持状态。
用例 |
|
|
|---|---|---|
在同一 TU 中定义和使用 |
支持 |
支持 |
在不同的 TU 中定义和使用 |
不支持 |
支持 |
在 -fno-gpu-rdc 模式下，编译器仅基于同一 TU 内存在的函数来计算内核的资源使用情况。此模式不支持使用在不同 TU 中定义的函数指针，因为这可能导致资源使用情况计算错误，从而导致未定义的行为。
另一方面，-fgpu-rdc 模式允许在不同的 TU 中定义和使用函数指针，因为资源使用情况计算可以适应来自不同 TU 的函数。
虚拟函数支持¶
在 Clang 与 HIP 中，在设备或主机代码中调用对象虚拟函数的支持取决于对象的构造位置。
**在设备代码中构造**：如果对象是在具有相同架构的卸载设备的设备代码中构造的，则可以在特定卸载设备的设备代码中调用对象的虚拟函数。
**在主机代码中构造**：如果对象是在主机代码中构造的，则可以在主机代码中调用对象的虚拟函数。
在其他情况下，不允许调用虚拟函数。
解释¶
在设备端构造的对象包含一个指向设备端虚拟函数表的指针，该指针在主机代码中不可访问，反之亦然。因此，尝试从与对象构造位置不同的上下文调用虚拟函数将被禁止，因为无法访问相应的虚拟表。不同架构的卸载设备的虚拟函数表不同，因此尝试从与对象构造位置不同的架构的卸载设备调用虚拟函数也会被禁止。
示例用法¶
class Base {
public:
__device__ virtual void virtualFunction() {
// Base virtual function implementation
}
};
class Derived : public Base {
public:
__device__ void virtualFunction() override {
// Derived virtual function implementation
}
};
__global__ void kernel() {
Derived obj;
Base* basePtr = &obj;
basePtr->virtualFunction(); // Allowed since obj is constructed in device code
}
HIPAMD 工具链上的 SPIR-V 支持¶
HIPAMD 工具链支持针对 AMDGCN 风格的 SPIR-V。ROCm 和 HIPAMD 工具链中对 SPIR-V 的支持正在积极开发中。
编译过程¶
当编译旨在利用 SPIR-V 的 HIP 程序时，该过程与传统的编译流程有所不同。
使用 --offload-arch=amdgcnspirv¶
**目标三元组**：
--offload-arch=amdgcnspirv标志指示编译器使用目标三元组spirv64-amd-amdhsa。这种方法会生成通用的 AMDGCN SPIR-V，它保留了特定于架构的元素，而不会对其进行硬编码，从而允许在运行时生成最佳的特定于目标的代码，此时具体目标已知。**LLVM IR 翻译**：程序被编译为 LLVM 中间表示（IR），随后被翻译成 SPIR-V。将来，此翻译步骤将被直接通过 SPIR-V 后端发出 SPIR-V 所取代。
**Clang 卸载捆绑器**：生成的 SPIR-V 被嵌入到 Clang 卸载捆绑器中，捆绑 ID 为
hip-spirv64-amd-amdhsa--amdgcnspirv。
与正常的 --offload-arch 混合¶
目前不支持通过 –offload-arch 混合使用 ``amdgcnspirv`` 和具体的 ``gfx###`` 目标；此限制是暂时的，将在未来版本中删除。
特定于架构的宏¶
当针对 SPIR-V 时，没有定义任何特定于架构的 AMDGPU 宏。将来会添加一个替代的、更灵活的机制来启用按目标/按功能的代码选择。