Transiting from CUDA to HIP （二）

一、Identifying Architecture Features

1. HIP_ARCH Defines

在 CUDA 编程中，__CUDA_ARCH__ 是一个预定义的宏，用于指示当前编译的代码所针对的 NVIDIA GPU 的计算能力（Compute Capability）。开发者可以使用这个宏来编写条件代码，以便在不同架构的 GPU 上支持不同的特性。例如，如果代码需要使用双精度浮点数（doubles），它可能会检查 __CUDA_ARCH__ 是否大于或等于 130，因为这是支持双精度浮点数的最小计算能力。

#if (__CUDA_ARCH__ >= 130) 
// doubles are supported 

然而，当使用 AMD GPU 和 ROCm 平台时，这种基于 __CUDA_ARCH__ 的检查就不再适用，因为 AMD GPU 有不同的架构和特性集。为了解决这个问题，HIP（Heterogeneous-compute Interface for Portability）提供了一套宏定义，使得开发者可以编写可移植的代码，这些代码可以在 NVIDIA 的 CUDA 平台和 AMD 的 ROCm 平台上运行。

HIP 提供的宏定义如下：

• __HIP_ARCH_HAS_DOUBLES__：如果当前的 GPU 支持双精度浮点数，则定义为真。

在编写 HIP 代码时，您应该使用这些宏来代替直接比较 __CUDA_ARCH__ 的值。这样做的好处是您的代码可以在不同的硬件架构上移植，而不需要为每种架构编写特定的条件代码

#if __HIP_ARCH_HAS_DOUBLES__
// 双精度浮点数在当前架构上是支持的
#endif

对于主机代码（即在 CPU 上运行的代码），所有 __HIP_ARCH__ 相关的宏都会被定义为 0。这意味着这些宏只应该在设备代码（即在 GPU 上运行的代码）中使用。

2. Device-Architecture Properties

在 CUDA 和 HIP 编程中，查询设备属性是一种常见的做法，用于确定特定 GPU 设备支持哪些特性。这对于编写可在不同 GPU 架构上运行的可移植代码非常重要。

在 HIP 中，hipGetDeviceProperties 函数用于获取设备的属性，这些属性描述了设备的能力和特性。设备属性结构体 hipDeviceProp_t 包含了多个字段，可以用来检查设备是否支持特定的功能。

以下是如何使用 hipGetDeviceProperties 来查询设备属性的示例：

hipDeviceProp_t deviceProp;
hipGetDeviceProperties(&deviceProp, device);

// 检查设备是否支持共享 int32 原子操作
if (deviceProp.arch.hasSharedInt32Atomics) {
    // 执行支持共享 int32 原子操作的代码
}

在这个例子中，deviceProp.arch.hasSharedInt32Atomics 是一个布尔字段，如果设备支持共享内存中的 32 位整数原子操作，则该字段为真。

与直接测试 major 和 minor 字段相比，使用 hipGetDeviceProperties 返回的结构体中的特定功能标志字段是一种更可移植的方法。直接测试 major 和 minor 字段通常与特定的 GPU 架构相关联，这限制了代码的可移植性：

// 不可移植的代码示例
if ((deviceProp.major == 1 && deviceProp.minor < 2)) {
    // 针对特定架构的代码
}

在编写 HIP 代码时，应该避免直接比较 major 和 minor 字段，而是使用 hipGetDeviceProperties 返回的属性来检查设备的功能。这样可以确保您的代码能够在不同的 GPU 设备上正常运行，无论它们是由 NVIDIA、AMD 还是其他制造商生产的。

3. Table of Architecture Properties

4. Finding HIP

在 Makefile 中，这段代码是一个条件赋值的示例，它使用了 ?= 运算符来为变量 HIP_PATH 提供一个默认值，但仅当该变量尚未设置时。这是一种常见的做法，用于在环境变量未预先定义的情况下提供一个默认值。

这里的 HIP_PATH 变量通常用于指定 HIP（Heterogeneous-compute Interface for Portability）的安装路径，而 hipconfig 是一个命令行工具，用于获取 HIP 相关的配置信息。

HIP_PATH ?= $(shell hipconfig --path)

• HIP_PATH：这是 Makefile 中的一个变量，用于存储 HIP 的安装路径。
• ?=：这是 Makefile 的条件赋值运算符。如果 HIP_PATH 变量在之前的 Makefile 或环境变量中没有被赋值，那么 ?= 运算符会将右侧的值赋给 HIP_PATH。
• $(shell hipconfig --path)：这是一个 shell 函数，它会执行 hipconfig --path 命令，并将其输出（即 HIP 的安装路径）赋值给 HIP_PATH 变量。hipconfig 工具通常用于获取 HIP 相关的配置信息，而 --path 选项用于获取 HIP 安装的根目录路径。

如果 HIP_PATH 已经在环境变量中设置或者在 Makefile 的其他地方被赋值，那么 $(shell hipconfig --path) 的结果不会被用来覆盖它。这样，用户就可以在需要时覆盖默认值。

5. Identifying HIP Runtime

在 HIP 的上下文中，"runtime" 指的是支持 HIP 代码执行的底层运行时环境。HIP 可以依赖于不同的运行时环境，具体取决于目标硬件平台：

1. ROCclr (Radeon Open Compute common language runtime)：

   • 在 AMD 平台上，HIP 使用 ROCclr 作为其运行时。ROCclr 是一个虚拟设备接口，它允许 HIP 运行时与不同的后端进行交互，包括 AMD 的 GPU。这使得 HIP 能够在 Linux 和 Windows 上运行，而无需对代码进行大量修改。

2. CUDA：

   • 在 NVIDIA 平台上，HIP 可以利用 CUDA 作为其运行时。在这种情况下，HIP 充当 CUDA API 之上的一层薄封装，使得原本为 CUDA 编写的代码可以几乎不做修改地在 HIP 上运行。

当在非 AMD 平台上使用 HIP 时，HIP 运行时会检查 CUDA 是否可用。如果检测到 CUDA，HIP 会设置 HIP_PLATFORM 环境变量为 NVIDIA，并使用 CUDA 的路径来编译和运行 HIP 代码。这允许开发者在 NVIDIA GPU 上利用 HIP 编写的代码，而无需对代码进行平台特定的修改。

6. Compiler Options

HIPcc 是 AMD 提供的一个编译器驱动程序，它是一个用于 HIP 应用程序的便携式编译器接口。根据目标系统，HIPcc 会调用 nvcc（NVIDIA 的 CUDA 编译器）或者 HIP-Clang（基于 LLVM 的编译器），并将所有必需的包含文件和库选项传递给目标编译器。这意味着 HIPcc 能够根据运行它的平台自动选择合适的编译器，并且设置正确的编译选项。

Compiler Options Supported on AMD Platforms

1. --amdgpu-target=：[已弃用] 此选项已被 --offload-arch= 替代。用于为目标 GPU 生成代码。支持的目标包括 gfx701, gfx801, gfx802, gfx803, gfx900, gfx906, gfx908, gfx1010, gfx1011, gfx1012, gfx1030, gfx1031。此选项可以在同一个命令行中多次出现，以生成支持多个目标的“胖二进制”（fat binary）。

2. --fgpu-rdc：生成可重定位的设备代码，允许内核或设备函数调用不同翻译单元中的设备函数。

3. -ggdb：等同于 -g 选项，并针对 GDB 进行调整。当使用 ROCm 的 GDB 调试 GPU 代码时，建议使用此选项。

4. --gpu-max-threads-per-block=：生成代码以支持每个块指定数量的线程。

5. -O：指定优化级别。

6. -offload-arch=：指定 AMD GPU 目标。这个选项用于替换 --amdgpu-target=。更多信息可以在 Clang 文档中找到，特别是关于目标 ID 的部分。

7. -save-temps：保存编译器生成的中间文件。

8. -show：显示编译步骤。

Option for specifying GPU processor

在 AMD 的 HIP (Heterogeneous-compute Interface for Portability) 编程指南中，--offload-arch=X 选项用于指定目标 GPU 的处理器或架构。这个选项告诉 hipcc 编译器为目标 AMD GPU 架构生成代码。这种指定方式有助于确保您的 HIP 应用程序能够在特定的硬件上运行。

示例：

hipcc --offload-arch=gfx908 my_program.cpp

7. Linking Issues

在使用 HIP（异构计算接口以提高可移植性）进行编程时，hipcc 是推荐的编译器驱动程序，因为它能够自动处理与 HIP 相关的库链接，以及管理 GPU 对象。

• hipcc 默认会在链接命令中添加 -lm 选项，这表示链接数学库（libm）。这个库包含了许多常用的数学函数，可能在您的应用程序中使用。
• 如果您需要添加额外的链接选项或库，可以在 hipcc 命令中指定它们。hipcc 会将这些选项传递给底层编译器。

二、Linking Code with Other Compilers

CUDA 代码通常使用 nvcc（NVIDIA CUDA 编译器）来处理加速器代码（定义和启动内核，通常在 .cu 或 .cuh 文件中定义）。同时，它也使用标准编译器（如 g++）来编译应用程序的其余部分。nvcc 是一个预处理器，它使用标准主机编译器（如 gcc）来生成主机代码。使用这个工具编译的代码只能使用 nvcc 和主机编译器都支持的语言特性的交集。在某些情况下，您需要确保主机编译器的数据类型和对齐方式与设备编译器的完全相同。只有一些主机编译器是支持的——例如，最近的 nvcc 版本缺乏对 Clang 主机编译器的支持。

相比之下，HIP-Clang 使用相同的基于 Clang 的编译器生成设备代码和主机代码。代码使用与 gcc 相同的 API，这允许由不同的 gcc 兼容编译器生成的代码相互链接。例如，使用 HIP-Clang 编译的代码可以与使用“标准”编译器（如 gcc、ICC 和 Clang）编译的代码链接。需要确保所有编译器使用相同的标准 C++ 头文件和库格式。

1. libc++ and libstdc++

hipcc 默认链接到 libstdc++ 库，这样做是为了在 g++ 和 HIP 之间提供更好的兼容性。libstdc++ 是 GNU 编译器集合（GCC）的 C++ 标准库实现，它被许多编译器支持，包括 g++。

如果您向 hipcc 传递 --stdlib=libc++ 选项，hipcc 将使用 libc++ 库。通常，libc++ 提供了更广泛的 C++ 特性集，而 libstdc++ 被更多编译器支持，尤其是 g++。

当交叉链接 C++ 代码时，任何使用 C++ 标准库中的类型（包括 std::string、std::vector 和其他容器）的 C++ 函数都必须使用相同的标准库实现。这包括以下情况：

• 在 HIP-Clang 中定义的函数或内核，从标准编译器调用。
• 在标准编译器中定义的函数，从 HIP-Clang 调用。
• 具有这些接口的应用程序应使用默认的 libstdc++ 链接。

如果您的应用程序完全用 hipcc 编译，并且从 libc++ 中受益于 libstdc++ 不支持的高级 C++ 特性，并且不需要与 nvcc 可移植性，您可以选择使用 libc++。

2. HIP Headers (hip_runtime.h, hip_runtime_api.h)

在 HIP (Heterogeneous-compute Interface for Portability) 编程中，hip_runtime.h 和 hip_runtime_api.h 是两个重要的头文件，它们定义了编译 HIP 程序所需的类型、函数和枚举：

1. hip_runtime_api.h：

   • 这个头文件定义了所有 HIP 运行时 API（例如 hipMalloc）以及调用它们所需的类型。
   • 如果一个源文件仅调用 HIP API，但不定义或启动任何内核，它可以包含 hip_runtime_api.h。
   • hip_runtime_api.h 不使用任何自定义的 hc（HIP Compiler）语言特性，因此可以用标准的 C++ 编译器编译。

2. hip_runtime.h：

   • 这个头文件被 hip_runtime_api.h 包含。它进一步提供了创建和启动内核所需的类型和定义。
   • 它可以用标准 C++ 编译器编译，但会暴露可用函数的一个子集。

与 CUDA 相比，这两个文件在 HIP 中的内容略有不同。在某些情况下，您可能需要将已经转换为 HIP 的代码修改为包含更丰富的 hip_runtime.h，而不是 hip_runtime_api.h。

使用建议：

• 如果您的代码只调用 HIP API 而不涉及内核定义或启动，使用 hip_runtime_api.h。
• 如果您的代码需要定义和启动内核，那么应该使用 hip_runtime.h，因为它提供了更全面的类型和定义支持。

// 使用 hip_runtime_api.h
#include <hip_runtime_api.h>

void allocateMemory() {
    float* devPtr;
    hipMalloc(&devPtr, sizeof(float) * 1024);
    // ...
}

// 使用 hip_runtime.h
#include <hip_runtime.h>

__global__ void myKernel(float* x, float* y) {
    // Kernel code
}

void launchKernel() {
    float* devPtrX, *devPtrY;
    hipMalloc(&devPtrX, sizeof(float) * 1024);
    hipMalloc(&devPtrY, sizeof(float) * 1024);
    
    hipLaunchKernel(myKernel, dim3(1, 1, 1), dim3(1024), 0, 0, devPtrX, devPtrY);
}

在第一个示例中，allocateMemory 函数仅调用 hipMalloc API，因此只包含 hip_runtime_api.h。 在第二个示例中，launchKernel 函数不仅调用 hipMalloc，还定义和启动了一个内核 myKernel，因此需要包含 hip_runtime.h。

3. Using a Standard C++ Compiler

在 HIP 编程环境中，hip_runtime_api.h 是一个关键的头文件，它包含了 HIP 运行时 API 的声明。这个头文件可以使用标准的 C 或 C++ 编译器（例如 gcc 或 ICC）来编译。为了正确编译包含 HIP 代码的文件，您需要确保编译器能够找到 HIP 的头文件，并且定义了正确的宏。

使用 hipconfig 获取编译选项

hipconfig 是一个工具，它提供了获取 HIP 编译和链接配置的命令行接口。您可以使用 hipconfig 来获取必要的编译器标志和包含路径。例如：

hipconfig --cxx_config -D__HIP_PLATFORM_AMD__ -I/home/user1/hip/include

这个命令会输出适合您环境的编译器标志和包含路径。

在 Makefile 中使用 hipconfig

您可以在 Makefile 中使用 hipconfig 来自动获取并设置编译器标志。以下是一个示例：

CPPFLAGS += $(shell $(HIP_PATH)/bin/hipconfig --cpp_config)

这行代码会将 hipconfig 返回的配置添加到 CPPFLAGS 变量中，这样在编译时就会包含正确的头文件路径和宏定义。

包含必要的 HIP 头文件

与 CUDA 不同，HIP 不会自动包含默认的头文件。因此，所有调用 HIP 运行时 API 或定义 HIP 内核的文件都必须显式包含相应的 HIP 头文件。如果编译过程中报告找不到必要的 API（例如，“错误：标识符 ‘hipSetDevice’ 未定义”），请确保文件包含了 hip_runtime.h（或在适当的情况下包含 hip_runtime_api.h）。

hipify-perl 脚本

hipify-perl 是一个用于将 CUDA 代码转换为 HIP 代码的脚本。它可以自动将 cuda_runtime.h 转换为 hip_runtime.h，并将 cuda_runtime_api.h 转换为 hip_runtime_api.h。但是，这个脚本可能会错过嵌套的头文件或宏定义，因此在使用后，您可能需要手动检查和调整代码。

cuda.h

在 HIP (Heterogeneous-compute Interface for Portability) 环境中，为了提高与 CUDA 代码的兼容性，HIP-Clang 提供了一个空的 cuda.h 文件。这是一个特殊的兼容性措施，用于处理那些包含了 cuda.h 但不实际使用其中任何函数的现有 CUDA 程序。

4. Choosing HIP File Extensions

在处理 HIP (Heterogeneous-compute Interface for Portability) 项目时，选择合适的文件扩展名可以帮助您区分不同类型的源代码文件，并确保构建系统正确地处理它们。以下是关于 HIP 文件扩展名选择的一些建议：

保留现有的 CUDA 文件扩展名

• .cu 和 .cuh：这些是 CUDA 项目中常用的文件扩展名，分别用于表示源文件和头文件。如果您正在将现有的 CUDA 项目快速迁移到 HIP，保留这些扩展名可能更简单，因为它减少了需要更改目录中文件名和 #include 语句的工作量。

推荐新的 HIP 文件扩展名

• .hip.cpp：对于源文件，推荐使用 .hip.cpp 扩展名。这表明该文件包含 HIP 代码，并且应该通过 hipcc 编译器进行编译。
• .hip.h 或 .hip.hpp：对于头文件，推荐使用 .hip.h 或 .hip.hpp 扩展名。这有助于区分 HIP 专用的头文件，并为构建工具提供明确的指示，以便在适当的时候调用 hipcc。

如果您正在创建一个新的 HIP 项目，您的文件结构可能如下所示：

project/
│
├── src/
│   ├── main.hip.cpp
│   └── vector_add.hip.cpp
│
├── include/
│   ├── vector_add.hip.h
│   └── utility.hip.hpp
│
└── Makefile

在 Makefile 中，您可以添加规则来处理 .hip.cpp 和 .hip.h 文件：

HIPCC = hipcc

# 编译 HIP 源文件
%.o: %.hip.cpp
	$(HIPCC) $(CPPFLAGS) $(CFLAGS) -c $< -o $@

# 包含目录
INCLUDES += -Iinclude

# 默认目标
all: main

main: main.o vector_add.o
	$(HIPCC) $(CFLAGS) $^ -o $@

# 包含依赖
-include $(DEPENDENCIES)

clean:
	rm -f *.o main

通过这种方式，您可以确保构建系统能够正确地识别和编译 HIP 代码，同时保持项目的组织和可维护性。