Transiting from CUDA to HIP（三）

青禾子的夏

于 2024-09-04 18:10:16 发布

阅读量1.2k

点赞数 20

文章标签： HIP 异构计算 Rocm 开发语言

本文链接：https://blog.csdn.net/weixin_44500921/article/details/141895073

版权

一、Workarounds

1. memcpyToSymbol

在 HIP (Heterogeneous-compute Interface for Portability) 中，hipMemcpyToSymbol 函数用于将数据从主机内存复制到设备上的全局内存或常量内存中，这样可以在设备端的内核中访问这些数据。这个功能特别有用，因为它允许在主机端定义数据符号，并在设备端的内核中使用这些符号。

#include<hip/hip_runtime.h> 
#include<hip/hip_runtime_api.h> 
#include<iostream> 
#define HIP_ASSERT(status) \ 
    assert(status == hipSuccess) 
#define LEN 512 
#define SIZE 2048 
__constant__ int Value[LEN]; // 定义了一个设备端的常量内存中的数组
__global__ void Get(hipLaunchParm lp, int *Ad) 
{ 
    int tid =threadIdx.x + blockIdx.x *blockDim.x; 
    Ad[tid] = Value[tid]; 
} 
int main() 
{ 
    int *A, *B, *Ad; 
    A = new int[LEN]; 
    B = new int[LEN]; 
    for(unsigned i=0;i<LEN;i++) 
    { 
        A[i] = -1*i; 
        B[i] = 0; 
    } 
    HIP_ASSERT(hipMalloc((void**)&Ad, SIZE)); 
    HIP_ASSERT(hipMemcpyToSymbol(HIP_SYMBOL(Value), A, SIZE, 0, hipMemcpyHostToDevice)); // 将主机端的 A 数组的数据复制到设备端的 Value 符号
    hipLaunchKernel(Get, dim3(1,1,1), dim3(LEN,1,1), 0, 0, Ad); // 启动内核 Get，它读取 Value 数组并将结果写入 Ad 数组
    HIP_ASSERT(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost)); 
    for(unsigned i=0;i<LEN;i++) 
    { 
        assert(A[i] == B[i]); 
    } 
    std::cout<<"Passed"<<std::endl; 
}

note: #define HIP_SYMBOL(name)：这是一个宏，用于将符号名称封装为 HIP 需要的格式。在调用 hipMemcpyToSymbol、hipMemcpyFromSymbol、hipGetSymbolAddress 和 hipGetSymbolSize 时，需要使用这个宏

2. CU_POINTER_ATTRIBUTE_MEMORY_TYPE

在 HIP (Heterogeneous-compute Interface for Portability) 中，hipPointerGetAttributes API 用于获取指针的属性，包括内存类型。这个功能在确定指针是分配在设备内存还是主机内存时非常有用。

#include <hip/hip_runtime.h>
#include <iostream>

int main() {
    // 分配设备内存
    double *ptr;
    hipMalloc(reinterpret_cast<void**>(&ptr), sizeof(double));

    // 获取设备指针的属性
    hipPointerAttribute_t attr;
    hipPointerGetAttributes(&attr, ptr);
    if (attr.memoryType == hipMemoryTypeDevice) {
        std::cout << "ptr is allocated on device memory." << std::endl;
    } else {
        std::cout << "ptr is not allocated on device memory." << std::endl;
    }

    // 分配主机内存
    double* ptrHost;
    hipHostMalloc(&ptrHost, sizeof(double));

    // 获取主机指针的属性
    hipPointerGetAttributes(&attr, ptrHost);
    if (attr.memoryType == hipMemoryTypeHost) {
        std::cout << "ptrHost is allocated on host memory." << std::endl;
    } else {
        std::cout << "ptrHost is not allocated on host memory." << std::endl;
    }

    // 清理资源
    hipFree(ptr);
    hipFreeHost(ptrHost);

    return 0;
}

3. threadfence_system

在 CUDA 和一些其他并行计算环境中，__threadfence_system() 函数用于确保当前线程发出的所有全局内存写入操作在所有设备上都可见。这包括对映射的主机内存和对等内存（peer memory）的写入。这个函数通常用于同步操作，确保在继续执行后续操作之前，所有先前的写入都已经完成。

然而，在 HIP (Heterogeneous-compute Interface for Portability) 中，并没有直接提供与 CUDA 中 __threadfence_system() 相同的功能。HIP 设计为与多种后端兼容，包括 AMD 的 ROCm 平台和 NVIDIA 的 CUDA 平台，因此它可能不包括特定于某一平台的同步机制。

由于 HIP 不直接提供 threadfence_system 功能，用户可以采取以下替代措施：

设置环境变量：
- 可以通过设置环境变量 HSA_DISABLE_CACHE=1 来禁用 GPU L2 缓存。这将影响所有访问和所有内核，可能会对性能产生影响。
- 这种方法是一种全局设置，可能会对整个应用程序的性能产生负面影响，因为它改变了 GPU 缓存的行为。
使用其他同步机制：
- 在 HIP 中，可以使用其他同步原语，如 hipDeviceSynchronize() 或 hipStreamSynchronize()，来确保内存操作的可见性和顺序。
- 这些函数可以确保所有先前的设备工作都已经完成，从而在一定程度上模拟 threadfence_system 的行为。

以下是使用 hipDeviceSynchronize() 来确保设备内存写入对 CPU 可见的示例：

#include <hip/hip_runtime.h>

__global__ void kernel() {
    // 执行一些内存写入操作
}

int main() {
    hipLaunchKernel(kernel, dim3(1), dim3(1), 0, 0);
    hipDeviceSynchronize(); // 等待设备完成所有先前排队的工作
    return 0;
}

4. Textures and Cache Control

在 GPU 编程中，纹理（texture）通常用于两种目的：利用专用的纹理缓存来加速内存访问，或者使用纹理采样硬件来进行插值和边界处理。这些用途在不同的硬件架构上可能会有不同的表现和优化方式。

纹理缓存与纹理采样硬件

纹理缓存：一些程序使用纹理来访问专用的纹理缓存，这可以通过简单的点采样器实现，基本上只读取一个点的数据。
纹理采样硬件：另一些程序则利用采样器硬件来进行插值和合并多个样本，这通常用于需要高级纹理处理功能的情况。

AMD 硬件与纹理缓存

AMD 的硬件以及一些竞争对手的较新硬件通常具有统一的纹理/L1 缓存，这意味着它们不再拥有专用的纹理缓存。

NVIDIA 硬件与纹理缓存

NVIDIA 的硬件（通过 nvcc 路径编译的 CUDA 程序）通常将全局加载的数据缓存到 L2 缓存中，这可能使得一些程序能够从显式控制 L1 缓存内容中受益。

__ldg 指令

__ldg 指令（Load Global）在 CUDA 中用于显式地从全局内存中加载数据，并且可以用于控制 L1 缓存的内容。在 AMD 硬件上，由于所有数据已经同时加载到 L1 和 L2 缓存中，__ldg 指令实际上被视为无操作（no-op）。

HIP 中的纹理和 __ldg 指令

对于只需要从改进的缓存中受益的程序，建议使用 __ldg 指令。
使用纹理对象和引用 API 的程序在 HIP 上表现良好，因为 HIP 支持这些功能。

功能性可移植性建议

对于只需要利用纹理来改善缓存的程序，可以使用 __ldg 指令，尽管在 AMD 硬件上它可能不会提供额外的性能优势。
对于使用纹理对象和引用 API 的程序，可以直接在 HIP 上运行，因为 HIP 提供了对这些特性的支持。

在 HIP 中使用 __ldg 指令的示例：

#include <hip/hip_runtime.h>

__global__ void kernel(float *data, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) {
        float value = __ldg(data + idx);
        // 进行一些操作
    }
}

int main() {
    float *data, *d_data;
    hipMalloc(&d_data, sizeof(float) * SIZE);
    // 假设 data 已经初始化
    hipMemcpy(d_data, data, sizeof(float) * SIZE, hipMemcpyHostToDevice);
    
    hipLaunchKernel(kernel, dim3(1, 1, 1), dim3(256, 1, 1), 0, 0, d_data, SIZE);
    
    hipFree(d_data);
    return 0;
}

二、More Tips

1. HIP Logging

HIP (Heterogeneous-compute Interface for Portability) 提供了日志记录功能，可以帮助开发者调试和分析 HIP 应用程序的执行情况。在 AMD 平台上，您可以通过设置环境变量 AMD_LOG_LEVEL 来控制日志记录的详细程度。

AMD_LOG_LEVEL 环境变量用于指定日志记录的级别，它支持不同的日志级别，每个级别对应不同的信息详细程度。以下是一些常用的日志级别：

0：关闭日志记录。
1：仅记录错误信息。
2：记录错误信息和警告信息。
3：记录错误、警告以及一些额外的调试信息。
4：记录详细的调试信息，包括函数调用、参数值等。

如何设置 AMD_LOG_LEVEL

您可以通过在命令行中设置环境变量或在程序中设置环境变量来控制日志级别。以下是两种设置方法：

在命令行中设置：
- 在 Linux 或 macOS 上，您可以在终端中使用以下命令：
```
export AMD_LOG_LEVEL=3
```
- 在 Windows 上，您可以在命令提示符或 PowerShell 中使用以下命令：
```
set AMD_LOG_LEVEL=3
```

在程序中设置：

您可以在程序启动时设置环境变量，例如在 C++ 程序中：

#include <stdlib.h>
#include <string.h>
#include <hip/hip_runtime.h>

int main() {
    // 设置日志级别
    setenv("AMD_LOG_LEVEL", "3", 1);

    // 您的 HIP 程序代码
    hipInit(0);
    // ...
    hipDeviceReset();
    return 0;
}

查看 HIP 日志

设置好 AMD_LOG_LEVEL 后，当您运行 HIP 应用程序时，日志信息将根据指定的级别输出到标准输出或指定的日志文件中。您可以在应用程序的输出中查看这些日志信息，以便进行调试和分析。

注意事项

日志记录可能会影响应用程序的性能，尤其是在较高的日志级别时，因为它会增加额外的输出操作。
确保在发布版本中关闭或减少日志记录的详细程度，以避免性能损失和不必要的信息泄露。

2. Debugging hipcc

在使用 HIP (Heterogeneous-compute Interface for Portability) 编译器 hipcc 进行编译时，如果需要查看 hipcc 产生的详细命令，可以通过设置环境变量 HIPCC_VERBOSE 来实现。这个环境变量的设置会让 hipcc 在标准错误输出（stderr）中打印出它生成的 HIP-clang（或 nvcc）命令。

您可以在命令行中设置 HIPCC_VERBOSE 环境变量，或者在程序中设置。以下是两种设置方法：

在命令行中设置：
- 在 Linux、macOS 或 Windows 的命令提示符（CMD）中，您可以设置环境变量：
```
export HIPCC_VERBOSE=1
```
- 在 Windows 的 PowerShell 中，您可以设置环境变量：
```
$env:HIPCC_VERBOSE="1"
```

在程序中设置：

您可以在 C++ 程序中使用 setenv 函数来设置环境变量：

#include <cstdlib>

int main() {
    // 设置 HIPCC_VERBOSE 环境变量
    setenv("HIPCC_VERBOSE", "1", 1);

    // 您的 HIP 程序代码
    // ...

    return 0;
}

查看编译命令

设置 HIPCC_VERBOSE 环境变量后，当您运行 hipcc 编译命令时，它将在标准错误输出中打印出详细的编译命令。这对于调试编译过程中的问题非常有用，因为它可以帮助您了解 hipcc 背后的实际编译过程。

示例

假设您有一个简单的 HIP 程序 example.cpp，您可以这样编译它：

hipcc example.cpp -o example

在设置了 HIPCC_VERBOSE 环境变量后，您会看到类似以下的输出：

# 例如，输出可能包括：
hipcc: executing HIP-clang command: hipcc -ccbin /usr/bin/clang -x hip --cuda -fcuda-rdc -o example.cpp.o -c example.cpp
hipcc: executing HIP-clang command: hipcc -ccbin /usr/bin/clang -o example example.cpp.o

这些输出显示了 hipcc 调用的实际编译器命令，包括所有的选项和参数。

注意事项

设置 HIPCC_VERBOSE 可能会产生大量的输出，因此建议仅在需要调试时使用。
确保在调试完成后，如果不需要详细输出，可以取消设置该环境变量或将其值设置回 0。

3. Editor Highlighting

为了在文本编辑器中获得 HIP (Heterogeneous-compute Interface for Portability) 文件的语法高亮，通常需要安装或配置相应的语法高亮插件或模式。许多流行的编辑器，如 Vim 和 Gedit，允许用户添加自定义的语法高亮规则来支持特定语言或框架。

Vim

对于 Vim 编辑器，您可以在 utils/vim 目录中找到 HIP 的语法高亮文件。通常，这个文件会被命名为 hip.vim，并且您可以通过以下步骤将其集成到 Vim 中：

将 hip.vim 文件复制到您的 Vim 配置目录下的 syntax 子目录中。在 Linux 系统上，这通常是 ~/.vim/syntax/。
如果 Vim 找不到该文件，您可能需要创建 syntax 目录：
```
mkdir -p ~/.vim/syntax
```
复制文件：
```
cp utils/vim/hip.vim ~/.vim/syntax/
```
为了让 Vim 知道如何使用这个新语法文件，您可能需要在 ~/.vimrc 文件中添加以下行：
```
autocmd BufRead,BufNewFile *.hip set filetype=hip
```
重新启动 Vim 并打开一个 .hip 文件，您应该能够看到语法高亮效果。

Gedit

对于 Gedit 编辑器，您可以在 utils/gedit 目录中找到 HIP 的语法高亮文件。这个文件可能是一个 .gedit-syntax 文件，并且您可以通过以下步骤将其添加到 Gedit：

将 .gedit-syntax 文件复制到 Gedit 的语法文件目录中。这通常是 ~/.local/share/gedit/syntaxes/。
如果 Gedit 语法文件目录不存在，您需要创建它：
```
mkdir -p ~/.local/share/gedit/syntaxes
```

复制文件：

cp utils/gedit/hip.gedit-syntax ~/.local/share/gedit/syntaxes/

重新启动 Gedit。您可能需要刷新 Gedit 的插件或重启计算机，以便 Gedit 识别新的语法文件。
打开一个 .hip 文件，然后在 Gedit 的“Preferences”菜单中选择“Editor” > “Highlighting and Bracketing”，并为您的文件类型选择正确的语法高亮。

注意事项

确保您复制的语法高亮文件与您的编辑器版本和配置兼容。
如果您在编辑器中没有看到预期的语法高亮效果，请检查您的文件扩展名是否正确，以及是否在编辑器的设置中启用了语法高亮。
一些编辑器可能需要额外的插件或工具来支持自定义语法高亮。

三、HIP Porting Driver API

1. Porting CUDA Driver API

在将 CUDA 代码移植到 HIP 时，需要注意 CUDA 提供了两套 API：CUDA Driver API 和 CUDA Runtime API。这两套 API 在功能上有显著的重叠，但也存在一些差异。以下是一些关键点，以及如何将 CUDA Driver API 移植到 HIP：

CUDA Driver API 与 Runtime API 的重叠功能

事件（Events）：cuEventCreate (Driver API) 和 cudaEventCreate (Runtime API)
流（Streams）：cuStreamCreate (Driver API) 和 cudaStreamCreate (Runtime API)
内存管理：cuMemAlloc (Driver API) 和 cudaMalloc (Runtime API)
内存复制：cuMemcpy (Driver API) 和 cudaMemcpy (Runtime API)
错误处理：cuGetErrorString (Driver API) 和 cudaGetErrorString (Runtime API)

CUDA Driver API 的独特功能

模块加载：cuModuleLoad (Driver API)，用于加载 CUDA 模块，没有直接的 Runtime API 等价物。
上下文管理：cuCtx 系列函数 (Driver API)，用于管理 CUDA 上下文，而 Runtime API 中没有直接的等价物。

错误代码和编码约定

CUDA Driver API 和 Runtime API 使用不同的错误代码空间和编码约定。例如，Driver API 使用 CUDA_ERROR_INVALID_VALUE，而 Runtime API 使用 cudaErrorInvalidValue。

移植 CUDA Driver API 到 HIP

在 HIP 中，大多数 CUDA Runtime API 都有直接的对应项，但 HIP 目前不支持与 CUDA Driver API 相同的模块加载和上下文管理功能。因此，移植 CUDA Driver API 到 HIP 时，您可能需要进行以下调整：

替换 API 调用：将 CUDA Driver API 调用替换为 HIP 运行时 API 调用。例如，将 cuEventCreate 替换为 hipEventCreate。
调整错误处理：将 CUDA Driver API 的错误代码替换为 HIP 的错误代码。HIP 错误代码通常与 CUDA Runtime API 的错误代码相似，但可能有所不同。
移除或替换模块和上下文管理：如果您的代码使用了 cuModule 或 cuCtx 系列函数，您可能需要重构代码以避免使用这些功能，或者寻找 HIP 提供的替代方案。
使用 HIP 提供的替代功能：对于 CUDA Driver API 中不直接由 HIP 支持的功能，您可能需要使用 HIP 提供的其他功能来实现类似的结果。

以下是将 CUDA Driver API 移植到 HIP 的示例：

// CUDA Driver API
// cuEvent_t event;
// cuEventCreate(&event, 0);

// HIP Runtime API
hipEvent_t event;
hipEventCreate(&event, 0);

2. cuModule API

在 CUDA 编程中，有两种 API（应用程序编程接口）可以用来与 GPU 交互：

CUDA Runtime API：
- 这是最常用的 API，它提供了一组函数，可以用来分配内存、创建执行流（streams）、同步操作等。
- 当你使用 Runtime API 时，所有的 CUDA 内核（kernels）都是自动编译并加载的。这意味着当你的程序运行时，它会包含所有的内核代码，并且 CUDA 运行时会自动处理加载这些内核的细节。
- 例如，如果你有一个名为 myKernel 的 CUDA 内核，你可以直接使用 <<< >>> 语法来启动它，这是 CUDA Runtime API 提供的一种快捷方式。
CUDA Driver API：
- 这是一个更底层的 API，它提供了更多的控制，比如可以显式地加载和卸载 CUDA 内核。
- 这个 API 允许你从文件或内存中加载 CUDA 模块（类似于动态链接库），并且可以从这些模块中提取内核函数。
- 例如，如果你有一个编译好的 CUDA 模块文件（.mycu），你可以使用 Driver API 来加载这个模块，然后获取模块中的特定内核函数并执行它。

在 HIP 中，也有类似的功能，但是它们是分开的：

HIP Runtime API：
- 与 CUDA Runtime API 类似，它提供了一组函数来执行 GPU 操作。
HIP Module API：
- 这是 HIP 提供的类似于 CUDA Driver API 的功能，它允许你加载和控制 CUDA 模块。

当你需要更细粒度的控制，比如只加载特定的内核或者在运行时生成内核代码时，使用 HIP Module API 就非常有用。这在一些高级用例中很有帮助，比如当你使用一种新的加速器语言或者有特殊的编译流程时。

总结一下，如果你的程序只需要使用标准的 CUDA 内核，并且你不需要对加载过程进行特别的控制，那么使用 HIP Runtime API 就足够了。但如果你有更复杂的需求，比如需要动态加载内核或者处理特殊的模块，那么你可能需要使用 HIP Module API。

以下是如何在 HIP 中使用模块 API 的示例：

#include <hip/hip_runtime.h>

int main() {
    hipModule_t module;
    hipFunction_t kernel;
    void* kernelArgs[] = {&arg1, &arg2};
    size_t numArgs = sizeof(kernelArgs) / sizeof(void*);

    // 加载模块
    hipModuleLoad(&module, "kernel.hip");

    // 获取内核函数
    hipModuleGetFunction(&kernel, module, "kernel_name");

    // 启动内核
    hipModuleLaunchKernel(kernel, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, 0, kernelArgs, nullptr);

    // 卸载模块
    hipModuleUnload(module);

    return 0;
}

在这个示例中，hipModuleLoad 用于加载代码对象，hipModuleGetFunction 用于获取内核函数，hipModuleLaunchKernel 用于启动内核，最后 hipModuleUnload 用于卸载模块。

3. cuCtx API

CUDA Driver API

在 CUDA Driver API 中，有两个重要的概念：设备（Device）和上下文（Context）。

设备（Device）：就像你的电脑里的 GPU 硬件。每个设备都有一个编号，比如设备 0、设备 1 等。
上下文（Context）：可以把它想象成一个工作环境，这个环境是针对一个特定设备的。在这个环境中，你可以创建和管理与那个设备相关的任务，比如运行内核（Kernel）或数据传输。

在 CUDA Driver API 中，你可以为每个设备创建一个上下文，然后在这个上下文中做很多事情，比如启动内核或创建事件（Event）。这种方式给你很多控制权，但也稍微复杂一些。

CUDA Runtime API

CUDA Runtime API 简化了这个过程。它没有明确区分上下文和设备，而是让你直接在设备上工作。你不需要创建上下文，而是直接告诉 CUDA 你想在哪个设备上工作。

HIP

HIP 是为了在不同的硬件平台上（比如 NVIDIA 的 GPU 和 AMD 的 GPU）都能运行而设计的。它在很多方面模仿了 CUDA Runtime API 的简化方式。

hipSetDevice：这个函数让你选择一个设备来工作，就像 CUDA Runtime API 中的选择设备一样。
流（Streams）：在 HIP 中，你可以创建流来管理设备上的任务。流允许你同时运行多个任务，提高了效率。

为什么上下文（Context）变得不那么重要了？

随着技术的发展，大多数情况下，你不需要创建和管理多个上下文。你只需要选择一个设备，然后直接在这个设备上运行你的代码。这就是为什么 HIP 和最新的 CUDA 版本都推荐使用更简单的方法来选择和使用设备。

总结

在 CUDA Driver API 中，你可以为每个 GPU 创建一个上下文，这给了你很多控制权，但也更复杂。
CUDA Runtime API 和 HIP 简化了这个过程，让你直接在设备上工作，不需要创建上下文。
HIP 提供了一些与 CUDA Driver API 相似的函数来帮助移植代码，但这些函数被认为是过时的，因为它们不符合简化的编程模型。

示例

在 CUDA Driver API 中，您可以这样创建和管理上下文：

CUcontext context;
CUdevice device;
cuDeviceGet(&device, 0); // 获取第一个设备
cuCtxCreate(&context, 0, device); // 创建上下文

在 HIP 中，您可以使用 hipSetDevice 来选择设备：

hipSetDevice(0); // 选择第一个设备

注意事项

在移植 CUDA Driver API 代码到 HIP 时，您可能需要调整与上下文相关的代码，因为 HIP 使用不同的机制来管理设备。
如果您的应用程序需要在多个 GPU 之间切换，您应该使用 HIP 提供的 hipSetDevice 或流 API。
由于 HIP 中的 hipCtx API 已被标记为弃用，建议避免在新代码中使用这些 API。

4. HIP Module and Ctx APIs

在 HIP (Heterogeneous-compute Interface for Portability) 中，为了提供一个统一的编程模型并简化开发过程，而不是提供两个独立的 API 集，HIP 通过引入新的模块（Modules）和上下文（Ctx）控制 API 来扩展其功能。这样做的目的是为了让 HIP 更容易从 CUDA 代码移植，同时保持 API 的一致性和简洁性。

模块（Module）API

在 CUDA 和 HIP 中，模块 API 允许你加载已经编译好的 GPU 代码。这些代码通常是由专门的编译器（比如 NVCC 或 HIP-Clang）生成的，并且存储在文件中。你可以把这些代码想象成 GPU 可以执行的程序。

NVCC 生成的代码通常有两种格式：.cubin 或 .ptx。
HIP-Clang 生成的代码格式为 .hsaco。

HIP 提供了 hipModuleLoad 这样的函数，让你可以从这些文件中加载代码到 GPU 中。这就像是在电脑上安装一个程序，一旦安装好了，你就可以运行它了。

上下文（Ctx）API

上下文 API 允许你管理 GPU 的工作环境。你可以把它想象成不同的工作区，每个工作区都可以有它自己的设置和任务。

在 CUDA 的早期版本中，每个 GPU 设备可以有多个这样的工作区，但 HIP 简化了这个过程，使得每个设备通常只有一个工作区。

为什么使用模块和上下文 API

模块 API：当你想要加载一些特别的或者在运行时生成的 GPU 代码时，这个 API 很有用。比如，你可能有一个程序，它可以根据用户的输入动态生成 GPU 代码，然后你需要加载这些代码来执行。
上下文 API：当你的程序需要在多个 GPU 之间切换，或者你需要更精细地控制 GPU 的运行环境时，这个 API 很有用。

示例

以下是使用 hipModule API 加载和使用代码对象的示例：

hipModule_t module;
hipFunction_t function;
void* kernelParams[] = {&arg1, &arg2};
size_t numParams = sizeof(kernelParams) / sizeof(void*);

// 加载模块
hipModuleLoad(&module, "kernel.hsaco");

// 获取模块中的函数
hipModuleGetFunction(&function, module, "kernel_func");

// 启动内核
hipModuleLaunchKernel(function, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, 0, kernelParams, nullptr);

// 卸载模块
hipModuleUnload(module);

5. hipCtx API

HIP 的 hipCtx API 提供了一种管理 GPU 上下文（context）的方法。在 HIP 中，上下文是一个抽象概念，它代表了与特定 GPU 设备相关的执行环境。这个 API 作为现有设备函数的一层薄封装，允许你设置当前的上下文或者查询与上下文相关联的设备属性。

hipCtx API 的主要功能包括：

设置当前上下文：你可以使用 hipCtx API 来指定你的应用程序当前应该使用的 GPU 上下文。这会影响到随后的 HIP 调用，因为它们会在这个当前上下文中执行。
查询设备属性：你可以查询与特定上下文相关联的 GPU 设备的属性，比如设备名称、计算能力、内存大小等。
与现有 HIP API 的集成：hipCtx API 与 HIP 的其他 API 紧密集成。例如，当你创建一个流（stream）时，如果没有明确指定上下文，HIP 会使用当前上下文来创建流。

使用 hipCtx API 的示例：

hipCtx_t ctx;
hipDevice_t device;

// 获取设备
hipDeviceGet(&device, 0);

// 创建上下文
hipCtxCreate(&ctx, 0, device);

// 设置当前上下文
hipCtxSetCurrent(ctx);

// 在当前上下文中创建一个流
hipStream_t stream;
hipStreamCreate(&stream);

// ... 执行一些操作 ...

// 销毁流
hipStreamDestroy(stream);

// 卸载当前上下文
hipCtxDestroy(ctx);

在这个示例中，我们首先获取了一个设备，然后创建了一个上下文。通过调用 hipCtxSetCurrent，我们设置了当前的上下文。这意味着任何后续的 HIP 调用都会在这个上下文中执行。我们还展示了如何在当前上下文中创建一个流，并在完成后销毁流和上下文。

注意事项：

上下文管理：在多设备环境中，正确管理上下文是很重要的。确保在适当的时间创建和销毁上下文，以避免资源泄漏。
性能考虑：虽然 hipCtx API 提供了额外的控制，但频繁地切换上下文可能会影响应用程序的性能。因此，合理规划上下文的使用是必要的。
API 弃用：在 HIP 的某些版本中，hipCtx API 可能被标记为弃用。这意味着 HIP 开发者推荐使用其他机制（如 hipSetDevice）来管理设备。在编写新的应用程序时，应考虑这一点。

总的来说，hipCtx API 提供了一种灵活的方式来管理 GPU 上下文，使得在多设备环境中的编程变得更加方便。然而，它也要求开发者对上下文的生命周期和使用有清晰的理解。

6. hipify translation of CUDA Driver API

hipify 是一个工具，它可以帮助开发者将 CUDA 代码转换为 HIP 代码。这个工具主要针对 CUDA Driver API，将其转换为 HIP 的等效调用。这种转换对于希望在 AMD GPU 上运行原有 CUDA 应用程序的开发者来说非常有用。以下是 hipify 工具处理的一些关键方面：

转换 CUDA Driver API 到 HIP

流（Streams）：cuStreamCreate 转换为 hipStreamCreate
事件（Events）：cuEventCreate 转换为 hipEventCreate
模块（Modules）：cuModuleLoad 转换为 hipModuleLoad
设备（Devices）：cuDeviceGet 转换为 hipDeviceGet
内存管理：cuMemAlloc 转换为 hipMemAlloc
上下文（Context）：cuCtxCreate 转换为 hipCtxCreate
性能分析（Profiler）：cuProfilerStart 转换为 hipProfilerStart

错误代码转换

hipify 工具还会将 CUDA Driver API 的错误代码转换为 HIP 的错误代码。例如，CUDA 中的 CUDA_ERROR_INVALID_VALUE 转换为 HIP 中的 hipErrorInvalidValue。

内存复制 API

CUDA Driver API 在内存复制函数的名称中包含了内存传输方向（例如，cuMemcpyH2D 表示从主机到设备）。
HIP 提供了两种风格的内存复制 API：
- 与 CUDA 类似的，包含方向的函数，如 hipMemcpyH2D。
- 单一的内存复制函数 hipMemcpy，它通过参数指定方向，并支持“默认”方向，让运行时自动确定方向。

性能考虑

使用明确指定内存方向的 API（如 hipMemcpyH2D）可能在某些情况下比使用单一函数（如 hipMemcpy）更快，因为它避免了主机开销来检测不同的内存方向。

错误处理

HIP 定义了一个统一的错误代码空间，并使用驼峰命名法（camel-case）为所有错误命名（例如 hipErrorInvalidValue）。

使用 hipify

要使用 hipify 工具，你通常需要在命令行中运行它，并指定要转换的 CUDA 源文件。hipify 会生成新的 HIP 源文件，你可以编译这些文件来在支持 HIP 的平台上运行。

hipify-perl your_cuda_file.cu -o your_hip_file.cpp

这个命令会将 your_cuda_file.cu 转换为 your_hip_file.cpp。

注意事项

hipify 工具可能无法处理所有的 CUDA 代码，特别是那些使用了 CUDA 特定功能或第三方库的代码。在这种情况下，可能需要手动进行一些修改。
转换后的代码可能需要进一步的调整和优化，以确保在新平台上的最佳性能。

四、HIP-Clang Implementation Notes

1. .hip_fatbin

在 HIP (Heterogeneous-compute Interface for Portability) 编程环境中，.hip_fatbin 是一种特殊的二进制格式，它允许将针对不同 GPU 架构的设备代码组合到一个单一的可执行文件或共享对象中。这种格式的主要目的是为了简化跨不同 GPU 架构的应用程序的部署和分发。

如何生成 .hip_fatbin

编译设备代码：当你使用 hip-clang 编译器编译 HIP 代码时，它会为每个目标设备生成一个代码对象（例如，.hsaco 文件）。
捆绑代码对象：clang-offload-bundler 工具会将这些针对不同设备的代码对象捆绑到一起，形成一个所谓的“fat binary”。
嵌入到可执行文件：这个 fat binary 会被嵌入到可执行文件或共享对象的 .hip_fatbin 节中，作为一个全局符号 __hfip_fatbin。

.hip_fatbin 的优点

简化部署：开发者不需要为每个目标设备提供单独的二进制文件。用户可以在支持的任何设备上运行同一个可执行文件，因为所需的设备代码已经包含在内。
易于分发：分发包含 .hip_fatbin 的应用程序更加方便，因为只需要一个包含所有目标设备代码的二进制文件。
自动选择：在运行时，HIP 运行时环境会自动选择与当前设备相匹配的代码对象进行加载和执行。

示例

假设你有两个 HIP 内核，分别针对不同的 GPU 架构。你可以使用 hip-clang 编译它们，并使用 clang-offload-bundler 将它们捆绑到一个 fat binary 中：

hip-clang -c -o kernel1.hsaco kernel1.cpp --offload-arch=gfx803
hip-clang -c -o kernel2.hsaco kernel2.cpp --offload-arch=gfx900
clang-offload-bundler -type=o -targets=hip-amd-gfx803-amdgcn--hip-amd-gfx900-amdgcn -inputs=kernel1.hsaco,kernel2.hsaco -outputs=kernels.hip_fatbin

然后，你可以将这个 .hip_fatbin 文件链接到你的应用程序中：

hip-clang -o my_application my_application.cpp -L. -lkernels

在这个例子中，-lkernels 选项告诉编译器链接 .hip_fatbin 文件，其中 -l 前缀表示这是一个库文件，而 kernels 是 .hip_fatbin 文件的名称。

2. Initialization and Termination Functions

在 HIP (Heterogeneous-compute Interface for Portability) 编程中，当你使用 HIP-Clang 编译主机代码时，它会为每个翻译单元生成初始化和终止函数。这些函数负责设置和清理与 GPU 相关的资源。以下是这些函数的作用和它们如何工作的详细说明：

初始化函数

__hipRegisterFatBinary：这个函数用于注册嵌入到 ELF 文件中的 .hip_fatbin 节的 fat binary。这是必要的，因为 fat binary 包含了针对不同 GPU 架构的设备代码。
__hipRegisterFunction：这个函数用于注册设备上的内核函数，使得它们可以被主机代码调用。
__hipRegisterVar：这个函数用于注册设备端的全局变量，使得它们可以在主机代码和设备代码之间共享。

终止函数

__hipUnregisterFatBinary：这个函数用于注销在程序开始时注册的 fat binary。这通常在程序结束时或在设备代码不再需要时调用。

全局变量 __hfip_gpubin_handle

HIP-Clang 为每个主机翻译单元生成一个名为 __hfip_gpubin_handle 的全局变量。这个变量的类型是 void**，并且它的链接属性被设置为 linkonce，这意味着在最终的二进制文件中，只会有一个这样的变量。

确保只注册一次

每个初始化函数都会检查 __hfip_gpubin_handle 的值。如果它是 0，这意味着 fat binary 还没有被注册，那么初始化函数会调用 __hipRegisterFatBinary 来注册 fat binary，并将返回值保存到 __hfip_gpubin_handle 中。这样做是为了保证 fat binary 只被注册一次。
类似地，终止函数也会检查 __hfip_gpubin_handle 的值，并在适当的时候调用 __hipUnregisterFatBinary 来注销 fat binary。

示例

这是一个简化的示例，展示了初始化和终止函数可能如何工作：

// 全局变量，用于检查是否已经注册了 fat binary
void* __hip_gpubin_handle = nullptr;

// 初始化函数
__attribute__((constructor))
void hipInit() {
    if (__hip_gpubin_handle == nullptr) {
        __hip_gpubin_handle = __hipRegisterFatBinary(&__hip_fatbin);
        // 注册内核函数和设备全局变量
        __hipRegisterFunction(&myKernel, ...); // 省略了参数
        __hipRegisterVar(&myDeviceVar, ...);  // 省略了参数
    }
}

// 终止函数
__attribute__((destructor))
void hipFini() {
    if (__hip_gpubin_handle != nullptr) {
        __hipUnregisterFatBinary(__hip_gpubin_handle);
        __hip_gpubin_handle = nullptr;
    }
}

在这个示例中，hipInit 函数是一个构造函数，它在程序启动时自动调用。它检查 __hfip_gpubin_handle 是否为 nullptr，如果是，则注册 fat binary。hipFini 函数是一个析构函数，它在程序结束时自动调用，用于注销 fat binary。

注意事项

这些初始化和终止函数是由 HIP-Clang 自动生成的，你通常不需要手动编写它们。
确保你的程序在适当的时机调用这些函数，以避免资源泄漏或未定义行为。

3. Kernel Launching

在 HIP-Clang 中，内核启动可以通过以下几种方式：

使用 CUDA 风格的三连字符语法 <<< >>>，这是最直接的方法，但在 HIP 中通常不推荐使用。
使用 hipLaunchKernel 宏，这是一个可变参数宏，可以接受启动配置（如网格维度、块维度、流、动态共享内存大小）和内核参数。这个宏会根据平台展开成适当的内核启动语法。
使用 hipLaunchKernelGGL 宏，这是 hipLaunchKernel 的一个变体，提供了一种标准 C/C++ 宏的方式来启动内核，可以作为 <<< >>> 语法的替代。

当使用 hipLaunchKernel 或 hipLaunchKernelGGL 宏时，需要提供内核名称、网格维度、块维度、动态共享内存大小和流等参数，然后是内核参数。这些宏的使用有助于代码的可移植性，并且可以在不同的 GPU 架构上运行。

此外，HIP-Clang 还提供了一些内置变量和函数，如 hipThreadIdx_x、hipBlockIdx_x、hipBlockDim_x、hipGridDim_x 等，这些与 CUDA 中的相应变量类似，用于在内核中确定线程的位置。

在迁移 CUDA 代码到 HIP 平台时，可以使用 hipify 工具自动转换大部分代码。这个工具可以处理许多常见的 CUDA 模式，但可能需要手动调整一些特定的架构特性查询或 HIP 不支持的 CUDA 功能。

4. Address Spaces

HIP-Clang 定义了一个进程范围的地址空间，其中 CPU 和所有设备都从一个统一的地址池中分配地址。因此，地址可以在不同上下文中共享，与原始 CUDA 定义不同，新的上下文不会为设备创建新的地址空间。这意味着在 HIP-Clang 中，内存管理是统一的，允许在 CPU 和 GPU 之间更容易地共享数据，而不需要在它们之间复制数据。

5. Using hipModuleLaunchKernel

在 HIP 中，hipModuleLaunchKernel 函数是与 CUDA 中的 cuLaunchKernel 相对应的函数。它用于启动内核执行，并且接受与 cuLaunchKernel 相同的参数。这些参数包括内核的执行配置，如网格维度、块维度、动态共享内存大小、流以及内核参数。

hipModuleLaunchKernel 的使用方式如下：

hipFunction_t function; // 已获取的内核函数指针
dim3 gridDim;          // 网格维度
dim3 blockDim;         // 块维度
size_t sharedMemBytes; // 动态共享内存大小
hipStream_t stream;    // 流

// 内核参数
void* kernelParams[] = {
    &arg1,
    &arg2,
    // ...
};

// 启动内核
hipModuleLaunchKernel(function, gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y, blockDim.z, sharedMemBytes, stream, kernelParams, 0);

6. Additional Information

在使用 HIP-Clang 编译 HIP 程序时，它会在调用 HIP API 时创建一个主上下文（primary context）。这个主上下文是 HIP-Clang 用来管理 GPU 设备资源和执行操作的上下文环境。在纯粹的驱动 API 代码中，这意味着当没有其他上下文存在时，HIP-Clang 会自动创建并使用这个主上下文。

与此相对的是，使用 HIP/NVCC（NVIDIA 的 CUDA 编译器）时，如果没有明确创建上下文，那么上下文栈将为空。在这种情况下，开发者需要手动创建和管理 CUDA 上下文。

当使用 HIP-Clang 时，如果上下文栈为空，HIP-Clang 会将主上下文推入上下文栈中。这样做的目的是为了确保在执行 HIP API 调用时有一个有效的上下文环境。这种自动管理上下文的行为可能会导致在使用运行时 API 和驱动 API 混合编程时出现微妙的差异，因为上下文的创建和销毁可能不会完全按照开发者的预期进行。

例如，如果开发者在代码中混合使用了 HIP 运行时 API 和驱动 API，那么在某些情况下，HIP-Clang 可能会自动创建和销毁上下文，而开发者可能并不期望这种行为。这可能会影响程序的性能，因为上下文的创建和销毁是一个相对昂贵的操作。

因此，在使用 HIP-Clang 进行编程时，了解上下文是如何被创建和管理的是很重要的。这有助于开发者更好地控制程序的行为，尤其是在涉及到跨多个 GPU 设备和上下文进行操作时。如果需要更精细的控制，开发者可以使用 HIP 的上下文管理 API 来显式地创建、设置和销毁上下文，以确保程序的行为符合预期。

五、NVCC Implementation Notes

1. Interoperation between HIP and CUDA Driver

CUDA 应用程序可能希望将 CUDA 驱动代码与 HIP 代码混合使用。下表显示了类型等价性，以实现这种交互。

2. Compilation Options

hipModule_t 接口不支持 cuModuleLoadDataEx 函数，该函数用于控制 PTX 编译选项。HIP-Clang 不使用 PTX，也不支持这些编译选项。HIP-Clang 的代码对象总是包含完全编译好的 ISA，不需要在加载步骤中进行额外的编译。相应的 HIP 函数 hipModuleLoadDataEx 在 HIP-Clang 路径上的行为与 hipModuleLoadData 相同（不使用编译选项），在 NVCC 路径上则表现为 cuModuleLoadDataEx。例如，CUDA 代码如下：

CUmodule module;
void *imagePtr = ...;  // 以某种方式用代码对象填充数据指针
const int numOptions = 1;
CUJit_option options[numOptions];
void * optionValues[numOptions];
options[0] = CU_JIT_MAX_REGISTERS;
unsigned maxRegs = 15;
optionValues[0] = (void*)(&maxRegs);

cuModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
CUfunction k;
cuModuleGetFunction(&k, module, "myKernel");

HIP 代码如下：

hipModule_t module;
void *imagePtr = ...;  // 以某种方式用代码对象填充数据指针
const int numOptions = 1;
hipJitOption options[numOptions];
void * optionValues[numOptions];
options[0] = hipJitOptionMaxRegisters;
unsigned maxRegs = 15;
optionValues[0] = (void*)(&maxRegs);
// 在 HIP-Clang 路径上将调用 hipModuleLoadData(module, imagePtr)，不会使用 JIT 选项，并且
// 在 NVCC 路径上将调用 hipModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues)
hipModuleLoadDataEx(module, imagePtr, numOptions, options, optionValues);
hipFunction_t k;
hipModuleGetFunction(&k, module, "myKernel");

下面的例子展示了如何使用 hipModuleGetFunction：

#include<hip_runtime.h>
#include<hip_runtime_api.h>
#include<iostream>
#include<fstream>
#include<vector>
#define LEN 64
#define SIZE LEN<<2
#ifdef __HIP_PLATFORM_HCC__
#define fileName "vcpy_isa.co"
#endif
#ifdef __HIP_PLATFORM_NVCC__
#define fileName "vcpy_isa.ptx"
#endif
#define kernel_name "hello_world"
int main(){
    float *A, *B;
    hipDeviceptr_t Ad, Bd;
    A = new float[LEN];
    B = new float[LEN];
    for(uint32_t i=0;i<LEN;i++){
        A[i] = i*1.0f;
        B[i] = 0.0f;
        std::cout<<A[i] << " "<<B[i]<<std::endl;
    }

#ifdef __HIP_PLATFORM_NVCC__
    hipInit(0);
    hipDevice_t device;
    hipCtx_t context;
    hipDeviceGet(&device, 0);
    hipCtxCreate(&context, 0, device);
#endif
    hipMalloc((void**)&Ad, SIZE);
    hipMalloc((void**)&Bd, SIZE);
    hipMemcpyHtoD(Ad, A, SIZE);
    hipMemcpyHtoD(Bd, B, SIZE);
    hipModule_t Module;
    hipFunction_t Function;
    hipModuleLoad(&Module, fileName);
    hipModuleGetFunction(&Function, Module, kernel_name);
    std::vector<void*>argBuffer(2);
    memcpy(&argBuffer[0], &Ad, sizeof(void*));
    memcpy(&argBuffer[1], &Bd, sizeof(void*));
    size_t size = argBuffer.size()*sizeof(void*);
    void *config[] = {
      HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
      HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
      HIP_LAUNCH_PARAM_END
    };
    hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config);
    hipMemcpyDtoH(B, Bd, SIZE);
    for(uint32_t i=0;i<LEN;i++){
        std::cout<<A[i]<<" - "<<B[i]<<std::endl;
    }

#ifdef __HIP_PLATFORM_NVCC__
    hipCtxDetach(context);
#endif

    return 0;
}

3. HIP Module and Texture Driver API

HIP 支持纹理驱动 API，但纹理引用应该在主机作用域内声明。以下代码解释了如何在 HIP_PLATFORM_HCC 平台上使用纹理引用。

// 生成代码对象的代码
#include "hip/hip_runtime.h"
extern texture<float, 2, hipReadModeElementType> tex;

__global__ void tex2dKernel(hipLaunchParm lp, float* outputData,
                            int width,
                            int height)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    outputData[y * width + x] = tex2D(tex, x, y);
}
// 主机代码
texture<float, 2, hipReadModeElementType> tex;
void myFunc()
{
    // ...
    textureReference* texref;
    hipModuleGetTexRef(&texref, Module1, "tex");
    hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap);
    hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap);
    hipTexRefSetFilterMode(texref, hipFilterModePoint);
    hipTexRefSetFlags(texref, 0);
    hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1);
    hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT);
    // ...
}

在这段代码中，我们首先在设备代码中定义了一个二维纹理 tex，然后在内核函数 tex2dKernel 中使用该纹理。在主机代码中，我们声明了一个与设备纹理相对应的主机侧纹理引用 tex，并通过 hipModuleGetTexRef 函数获取纹理引用的指针 texref。接着，我们使用一系列函数来设置纹理的地址模式、过滤模式、标志、格式和数组。