memcpyToSymbol

函数用于将数据从主机内存复制到设备上的全局内存或常量内存中这样可以在设备端的内核中访问这些数据。

这个功能特别有用因为它允许在主机端定义数据符号并在设备端的内核中使用这些符号。

#includeiostream

HIP_ASSERT(hipMalloc((void**)Ad,

SIZE));

HIP_ASSERT(hipMemcpyToSymbol(HIP_SYMBOL(Value),

SIZE,

HIP_SYMBOL(name)这是一个宏用于将符号名称封装为

HIP

hipMemcpyToSymbol、hipMemcpyFromSymbol、hipGetSymbolAddress

hipGetSymbolSize

CU_POINTER_ATTRIBUTE_MEMORY_TYPE

HIP

用于获取指针的属性包括内存类型。

这个功能在确定指针是分配在设备内存还是主机内存时非常有用。

#include

*ptr;hipMalloc(reinterpret_castvoid**(ptr),

sizeof(double));//

attr;hipPointerGetAttributes(attr,

ptr);if

获取主机指针的属性hipPointerGetAttributes(attr,

ptrHost);if

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

threadfence_system

和一些其他并行计算环境中__threadfence_system()

函数用于确保当前线程发出的所有全局内存写入操作在所有设备上都可见。

这包括对映射的主机内存和对等内存peer

memory的写入。

这个函数通常用于同步操作确保在继续执行后续操作之前所有先前的写入都已经完成。

然而在

缓存。

这将影响所有访问和所有内核可能会对性能产生影响。

这种方法是一种全局设置可能会对整个应用程序的性能产生负面影响因为它改变了

GPU

hipStreamSynchronize()来确保内存操作的可见性和顺序。

这些函数可以确保所有先前的设备工作都已经完成从而在一定程度上模拟

threadfence_system

编程中纹理texture通常用于两种目的利用专用的纹理缓存来加速内存访问或者使用纹理采样硬件来进行插值和边界处理。

这些用途在不同的硬件架构上可能会有不同的表现和优化方式。

纹理缓存与纹理采样硬件

纹理缓存一些程序使用纹理来访问专用的纹理缓存这可以通过简单的点采样器实现基本上只读取一个点的数据。

纹理采样硬件另一些程序则利用采样器硬件来进行插值和合并多个样本这通常用于需要高级纹理处理功能的情况。

AMD

hipMemcpyHostToDevice);hipLaunchKernel(kernel,

dim3(1,

环境变量用于指定日志记录的级别它支持不同的日志级别每个级别对应不同的信息详细程度。

以下是一些常用的日志级别

0关闭日志记录。

1仅记录错误信息。

2记录错误信息和警告信息。

3记录错误、警告以及一些额外的调试信息。

4记录详细的调试信息包括函数调用、参数值等。

如何设置

您可以通过在命令行中设置环境变量或在程序中设置环境变量来控制日志级别。

以下是两种设置方法

在命令行中设置

应用程序时日志信息将根据指定的级别输出到标准输出或指定的日志文件中。

您可以在应用程序的输出中查看这些日志信息以便进行调试和分析。

注意事项

日志记录可能会影响应用程序的性能尤其是在较高的日志级别时因为它会增加额外的输出操作。

确保在发布版本中关闭或减少日志记录的详细程度以避免性能损失和不必要的信息泄露。

Debugging

编译命令时它将在标准错误输出中打印出详细的编译命令。

这对于调试编译过程中的问题非常有用因为它可以帮助您了解

hipcc

可能会产生大量的输出因此建议仅在需要调试时使用。

确保在调试完成后如果不需要详细输出可以取消设置该环境变量或将其值设置回

Editor

文件的语法高亮通常需要安装或配置相应的语法高亮插件或模式。

许多流行的编辑器如

Vim

Gedit允许用户添加自定义的语法高亮规则来支持特定语言或框架。

Vim

~/.local/share/gedit/syntaxes/。

Gedit

确保您复制的语法高亮文件与您的编辑器版本和配置兼容。

如果您在编辑器中没有看到预期的语法高亮效果请检查您的文件扩展名是否正确以及是否在编辑器的设置中启用了语法高亮。

一些编辑器可能需要额外的插件或工具来支持自定义语法高亮。

三、HIP

在功能上有显著的重叠但也存在一些差异。

以下是一些关键点以及如何将

CUDA

API它提供了一组函数可以用来分配内存、创建执行流streams、同步操作等。

当你使用

Runtime

内核kernels都是自动编译并加载的。

这意味着当你的程序运行时它会包含所有的内核代码并且

CUDA

模块类似于动态链接库并且可以从这些模块中提取内核函数。

例如如果你有一个编译好的

CUDA

当你需要更细粒度的控制比如只加载特定的内核或者在运行时生成内核代码时使用

HIP

就非常有用。

这在一些高级用例中很有帮助比如当你使用一种新的加速器语言或者有特殊的编译流程时。

CUDA

就足够了。

但如果你有更复杂的需求比如需要动态加载内核或者处理特殊的模块那么你可能需要使用

HIP

获取内核函数hipModuleGetFunction(kernel,

module,

启动内核hipModuleLaunchKernel(kernel,

gridDimX,

卸载模块hipModuleUnload(module);return

在这个示例中hipModuleLoad

等。

上下文Context可以把它想象成一个工作环境这个环境是针对一个特定设备的。

在这个环境中你可以创建和管理与那个设备相关的任务比如运行内核Kernel或数据传输。

CUDA

中你可以为每个设备创建一个上下文然后在这个上下文中做很多事情比如启动内核或创建事件Event。

这种方式给你很多控制权但也稍微复杂一些。

CUDA

简化了这个过程。

它没有明确区分上下文和设备而是让你直接在设备上工作。

你不需要创建上下文而是直接告诉

CUDA

中你可以创建流来管理设备上的任务。

流允许你同时运行多个任务提高了效率。

随着技术的发展大多数情况下你不需要创建和管理多个上下文。

你只需要选择一个设备然后直接在这个设备上运行你的代码。

这就是为什么

HIP

相似的函数来帮助移植代码但这些函数被认为是过时的因为它们不符合简化的编程模型。

CUDA

中为了提供一个统一的编程模型并简化开发过程而不是提供两个独立的

API

HIP-Clang生成的并且存储在文件中。

你可以把这些代码想象成

GPU

中。

这就像是在电脑上安装一个程序一旦安装好了你就可以运行它了。

上下文CtxAPI

的工作环境。

你可以把它想象成不同的工作区每个工作区都可以有它自己的设置和任务。

CUDA

hipModuleLaunchKernel(function,

gridDimX,

作为现有设备函数的一层薄封装允许你设置当前的上下文或者查询与上下文相关联的设备属性。

hipCtx

紧密集成。

例如当你创建一个流stream时如果没有明确指定上下文HIP

hipCtx

在这个示例中我们首先获取了一个设备然后创建了一个上下文。

通过调用

hipCtxSetCurrent我们设置了当前的上下文。

这意味着任何后续的

HIP

调用都会在这个上下文中执行。

我们还展示了如何在当前上下文中创建一个流并在完成后销毁流和上下文。

注意事项

上下文管理在多设备环境中正确管理上下文是很重要的。

确保在适当的时间创建和销毁上下文以避免资源泄漏。

性能考虑虽然

hipCtx

提供了额外的控制但频繁地切换上下文可能会影响应用程序的性能。

因此合理规划上下文的使用是必要的。

API

弃用在

hipSetDevice来管理设备。

在编写新的应用程序时应考虑这一点。

API

上下文使得在多设备环境中的编程变得更加方便。

然而它也要求开发者对上下文的生命周期和使用有清晰的理解。

hipify

hipStreamCreate事件EventscuEventCreate

转换为

hipEventCreate模块ModulescuModuleLoad

转换为

hipModuleLoad设备DevicescuDeviceGet

转换为

hipMemAlloc上下文ContextcuCtxCreate

转换为

hipCtxCreate性能分析ProfilercuProfilerStart

转换为

在内存复制函数的名称中包含了内存传输方向例如cuMemcpyH2D

API

hipMemcpy它通过参数指定方向并支持“默认”方向让运行时自动确定方向。

性能考虑

hipMemcpy更快因为它避免了主机开销来检测不同的内存方向。

错误处理

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

hipify

特定功能或第三方库的代码。

在这种情况下可能需要手动进行一些修改。

转换后的代码可能需要进一步的调整和优化以确保在新平台上的最佳性能。

四、HIP-Clang

架构的设备代码组合到一个单一的可执行文件或共享对象中。

这种格式的主要目的是为了简化跨不同

GPU

工具会将这些针对不同设备的代码对象捆绑到一起形成一个所谓的“fat

fat

简化部署开发者不需要为每个目标设备提供单独的二进制文件。

用户可以在支持的任何设备上运行同一个可执行文件因为所需的设备代码已经包含在内。

易于分发分发包含

.hip_fatbin

的应用程序更加方便因为只需要一个包含所有目标设备代码的二进制文件。

自动选择在运行时HIP

运行时环境会自动选择与当前设备相匹配的代码对象进行加载和执行。

假设你有两个

-targetship-amd-gfx803-amdgcn--hip-amd-gfx900-amdgcn

-inputskernel1.hsaco,kernel2.hsaco

.hip_fatbin

编译主机代码时它会为每个翻译单元生成初始化和终止函数。

这些函数负责设置和清理与

GPU

__hipRegisterFatBinary这个函数用于注册嵌入到

ELF

__hipRegisterFunction这个函数用于注册设备上的内核函数使得它们可以被主机代码调用。

__hipRegisterVar这个函数用于注册设备端的全局变量使得它们可以在主机代码和设备代码之间共享。

终止函数

__hipUnregisterFatBinary这个函数用于注销在程序开始时注册的

fat

linkonce这意味着在最终的二进制文件中只会有一个这样的变量。

确保只注册一次

__hipRegisterFatBinary(__hip_fatbin);//

注册内核函数和设备全局变量__hipRegisterFunction(myKernel,

...);

省略了参数__hipRegisterVar(myDeviceVar,

...);

{__hipUnregisterFatBinary(__hip_gpubin_handle);__hip_gpubin_handle

nullptr;}

自动生成的你通常不需要手动编写它们。

确保你的程序在适当的时机调用这些函数以避免资源泄漏或未定义行为。

Kernel

宏这是一个可变参数宏可以接受启动配置如网格维度、块维度、流、动态共享内存大小和内核参数。

这个宏会根据平台展开成适当的内核启动语法。

使用

hipLaunchKernelGGL

宏时需要提供内核名称、网格维度、块维度、动态共享内存大小和流等参数然后是内核参数。

这些宏的使用有助于代码的可移植性并且可以在不同的

GPU

hipThreadIdx_x、hipBlockIdx_x、hipBlockDim_x、hipGridDim_x

等这些与

和所有设备都从一个统一的地址池中分配地址。

因此地址可以在不同上下文中共享与原始

CUDA

相同的参数。

这些参数包括内核的执行配置如网格维度、块维度、动态共享内存大小、流以及内核参数。

hipFunction_t

hipModuleLaunchKernel(function,

gridDim.x,

编译器时如果没有明确创建上下文那么上下文栈将为空。

在这种情况下开发者需要手动创建和管理

CUDA

调用时有一个有效的上下文环境。

这种自动管理上下文的行为可能会导致在使用运行时

API

混合编程时出现微妙的差异因为上下文的创建和销毁可能不会完全按照开发者的预期进行。

HIP

可能会自动创建和销毁上下文而开发者可能并不期望这种行为。

这可能会影响程序的性能因为上下文的创建和销毁是一个相对昂贵的操作。

因此在使用

进行编程时了解上下文是如何被创建和管理的是很重要的。

这有助于开发者更好地控制程序的行为尤其是在涉及到跨多个

GPU

(void*)(maxRegs);cuModuleLoadDataEx(module,

imagePtr,

__HIP_PLATFORM_NVCC__hipInit(0);hipDevice_t

device;hipCtx_t

fileName);hipModuleGetFunction(Function,

Module,

kernel_name);std::vectorvoid*argBuffer(2);memcpy(argBuffer[0],

Ad,

sizeof(void*));memcpy(argBuffer[1],

Bd,

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,

LEN,

(void**)config);hipMemcpyDtoH(B,

Bd,

__HIP_PLATFORM_NVCC__hipCtxDetach(context);

#endifreturn

texref;hipModuleGetTexRef(texref,

Module1,

tex);hipTexRefSetAddressMode(texref,

hipAddressModeWrap);hipTexRefSetAddressMode(texref,

hipAddressModeWrap);hipTexRefSetFilterMode(texref,

hipFilterModePoint);hipTexRefSetFlags(texref,

HIP_AD_FORMAT_FLOAT,

中使用该纹理。

在主机代码中我们声明了一个与设备纹理相对应的主机侧纹理引用

hipModuleGetTexRef

texref。

接着我们使用一系列函数来设置纹理的地址模式、过滤模式、标志、格式和数组。