贵阳网站设计阳光创信好吗,lamp wordpress 404,现在搜什么关键词能搜到网站,专业的购物网站建设一、Workarounds
1. memcpyToSymbol
在 HIP (Heterogeneous-compute Interface for Portability) 中#xff0c;hipMemcpyToSymbol 函数用于将数据从主机内存复制到设备上的全局内存或常量内存中#xff0c;这样可以在设备端的内核中访问这些数据。这个功能特别有用#x…一、Workarounds
1. memcpyToSymbol
在 HIP (Heterogeneous-compute Interface for Portability) 中hipMemcpyToSymbol 函数用于将数据从主机内存复制到设备上的全局内存或常量内存中这样可以在设备端的内核中访问这些数据。这个功能特别有用因为它允许在主机端定义数据符号并在设备端的内核中使用这些符号。
#includehip/hip_runtime.h
#includehip/hip_runtime_api.h
#includeiostream
#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 i0;iLEN;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 i0;iLEN;i) { assert(A[i] B[i]); } std::coutPassedstd::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 iostreamint main() {// 分配设备内存double *ptr;hipMalloc(reinterpret_castvoid**(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_CACHE1 来禁用 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_LEVEL3 在 Windows 上您可以在命令提示符或 PowerShell 中使用以下命令 set AMD_LOG_LEVEL3 在程序中设置 您可以在程序启动时设置环境变量例如在 C 程序中 #include stdlib.h
#include string.h
#include hip/hip_runtime.hint 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_VERBOSE1 在 Windows 的 PowerShell 中您可以设置环境变量 $env:HIPCC_VERBOSE1 在程序中设置 您可以在 C 程序中使用 setenv 函数来设置环境变量 #include cstdlibint 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 filetypehip 重新启动 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 提供了两套 APICUDA Driver API 和 CUDA Runtime API。这两套 API 在功能上有显著的重叠但也存在一些差异。以下是一些关键点以及如何将 CUDA Driver API 移植到 HIP
CUDA Driver API 与 Runtime API 的重叠功能
事件EventscuEventCreate (Driver API) 和 cudaEventCreate (Runtime API)流StreamscuStreamCreate (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.hint 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 的一致性和简洁性。
模块ModuleAPI
在 CUDA 和 HIP 中模块 API 允许你加载已经编译好的 GPU 代码。这些代码通常是由专门的编译器比如 NVCC 或 HIP-Clang生成的并且存储在文件中。你可以把这些代码想象成 GPU 可以执行的程序。
NVCC 生成的代码通常有两种格式.cubin 或 .ptx。HIP-Clang 生成的代码格式为 .hsaco。
HIP 提供了 hipModuleLoad 这样的函数让你可以从这些文件中加载代码到 GPU 中。这就像是在电脑上安装一个程序一旦安装好了你就可以运行它了。
上下文CtxAPI
上下文 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
流StreamscuStreamCreate 转换为 hipStreamCreate事件EventscuEventCreate 转换为 hipEventCreate模块ModulescuModuleLoad 转换为 hipModuleLoad设备DevicescuDeviceGet 转换为 hipDeviceGet内存管理cuMemAlloc 转换为 hipMemAlloc上下文ContextcuCtxCreate 转换为 hipCtxCreate性能分析ProfilercuProfilerStart 转换为 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-archgfx803
hip-clang -c -o kernel2.hsaco kernel2.cpp --offload-archgfx900
clang-offload-bundler -typeo -targetship-amd-gfx803-amdgcn--hip-amd-gfx900-amdgcn -inputskernel1.hsaco,kernel2.hsaco -outputskernels.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/NVCCNVIDIA 的 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
#includehip_runtime.h
#includehip_runtime_api.h
#includeiostream
#includefstream
#includevector
#define LEN 64
#define SIZE LEN2
#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 i0;iLEN;i){A[i] i*1.0f;B[i] 0.0f;std::coutA[i] B[i]std::endl;}#ifdef __HIP_PLATFORM_NVCC__hipInit(0);hipDevice_t device;hipCtx_t context;hipDeviceGet(device, 0);hipCtxCreate(context, 0, device);
#endifhipMalloc((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::vectorvoid*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 i0;iLEN;i){std::coutA[i] - B[i]std::endl;}#ifdef __HIP_PLATFORM_NVCC__hipCtxDetach(context);
#endifreturn 0;
}
3. HIP Module and Texture Driver API
HIP 支持纹理驱动 API但纹理引用应该在主机作用域内声明。以下代码解释了如何在 HIP_PLATFORM_HCC 平台上使用纹理引用。
// 生成代码对象的代码
#include hip/hip_runtime.h
extern texturefloat, 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);
}
// 主机代码
texturefloat, 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。接着我们使用一系列函数来设置纹理的地址模式、过滤模式、标志、格式和数组。
