ROCm技术解析概述

3.2 ROCm技术解析

ROCm是第一个用于GPU计算的开源HPC/Hyperscale级平台，也是独立于编程语言的。将UNIX的选择哲学、极简主义和模块化软件开发引入GPU计算。新的ROCm基础允许为应用程序选择甚至开发工具和语言运行时。

1）[n1] ROCm是为规模而构建的；它支持通过RDMA进行服务器节点内外的多GPU计算通信。当驱动程序直接结合RDMA对等同步支持时，它还简化了堆栈。

ROCm具有丰富的系统运行时，具有大规模应用程序、编译器和语言运行时开发所需的关键功能。

1. 走向11：增强编程语言运行时基础

ROCr系统运行时是独立于语言的，并大量使用异构系统架构（HSA）运行时API。这种方法为执行编程语言提供了丰富的基础，如HCC C++和HIP、Khronos Group的OpenCL和Continuum的Anaconda Python。

重要功能包括以下内容：

1）多GPU粗粒度共享虚拟内存

2）进程并发和抢占

3）大内存分配

4）HSA信号和原子

5）用户模式队列和DMA

6）标准化加载器和代码对象格式

7）动态和离线编译支持

8）支持RDMA的点对点多GPU操作

9）事件探查器跟踪和事件收集API

10）系统管理API和工具

    ROCm支持苹果公司Metal，如图3-4所示。

 

图3-4 ROCm支持苹果公司Metal

2. 坚实的编译基础和语言支持

1）LLVM编译器基础

2）HCC C++和HIP用于应用程序可移植性

3）GCN组装器和拆卸器

ROCm所能实现的领域[n2] 是广阔而未知的。

3. 上游Linux内核中的ROCm支持

从ROCm 1.9.0开始，ROCm用户级软件与某些上游Linux内核中的AMD驱动程序兼容。因此，用户可以选择使用作为AMD ROCm存储库一部分的ROCK内核驱动程序，也可以选择使用上游驱动程序，只从AMD的ROCm存储库安装ROCm用户级实用程序。

---

 [n1]只有1？其他序号没了

 [n2]病句，表达不清

这些版本的上游Linux内核支持ROCm中的以下GPU：

1）4.17: Fiji, Polaris 10, Polaris 11

2）4.18: Fiji, Polaris 10, Polaris 11, Vega10

3）4.20: Fiji, Polaris 10, Polaris 11, Vega10, Vega 7nm

上游驱动程序可能有助于在与AMD存储库中，可用的内核驱动程序不兼容系统运行ROCm软件。对于可以选择使用AMD或上游驱动程序的用户，需要考虑各种权衡，见表3-1。

表3-1 对于可以选择使用AMD或上游驱动程序的用户，需要考虑各种权衡

	使用AMD的rock dkms软件包	使用上游内核驱动程序
正面的	更多GPU功能，并且它们更早启用	包括最新的Linux内核功能
	AMD在支持的发行版上进行了测试	可能适用于其他发行版和自定义内核
	支持的GPU已启用，无论内核版本如何
	包含最新的GPU固件
反面的	可能不适用于所有Linx发行版或版本	功能和硬件支持因内核版本而异
	4.18以上的内核目前不支持	将GPU对系统内存的使用限制为系统内存的3/8
	IPC和RDMA功能尚未启用
	AMD没有测试到与rock dkms软件包相同的级别
	不包括最新固件

4. 从AMD ROCm存储库安装

AMD目前为ROCm 2.4.x软件包托管Debian和RPM存储库。Debian存储库中的软件包已经过签名，以确保软件包的完整性。

5. ROCm二进制包结构

ROCm是一系列软件的集合，从驱动程序和运行时到库和开发人员工具。在AMD的软件包发行版中，这些软件项目作为单独的软件包提供。如果用户不想安装所有ROCm，允许他们只安装所需的软件包。默认情况下，这些软件包会将大部分ROCm软件安装到/opt/ROCm/中。

6. 小结一下

从高层次的角度来看，ROCm提供了一套丰富的工具，允许为应用程序选择最佳语言。

1）HCC（异构计算编译器）支持HC方言

2）HIP是一个运行时库，它位于HCC之上（对于AMD ROCm平台；对于Nvidia，它使用NVCC编译器）

3）以下内容将很快为GCN ISA提供本机编译器支持：

①OpenCL 1.2+

②Anaconda（Python）与Numba

所有这些都是开源项目，所以可以从语言到硬件都采用完全开放的堆栈。

比较不同计算API的语法，见表3-2。

表3-2 比较不同计算API的语法

Term	CUDA	HIP	HC	C++AMP	OpenCL
设备	int deviceId	int deviceId	hc::accelerator	concurrency:: accelerator	cl_device
队列	cudaStream_t	hipStream_t	hc:: accelerator_view	concurrency:: accelerator_view	cl_command_queue
事件	cudaEvent_t	hipEvent_t	hc:: completion_future	concurrency:: completion_future	cl_event
内存	void *	void *	void*; hc::array; hc::array_view	concurrency::array; concurrency::array_view	cl_mem
	网格 块 线程 变形	网格   块 线程   变形	程度   小块 线程   波前	程度 小块 线程   N/A	NDRange work-group work-item sub-group
线程索引	threadIdx.x	hipThreadIdx_x	t_idx.local[0]	t_idx.local[0]	get_local_id(0)
块索引	blockIdx.x	hipBlockIdx_x	t_idx.tile[0]	t_idx.tile[0]	get_group_id(0)
块维数	blockDim.x	hipBlockDim_x	t_ext.tile_dim[0]	t_idx.tile_dim0	get_local_size(0)
Grid-dim	gridDim.x	hipGridDim_x	t_ext[0]	t_ext[0]	get_global_size(0)
设备功能	__device__	__device__	[[hc]] (在许多情况下自动检测)	restrict(amp)	隐含在设备编译中
主机功能	__host_ (default)	__host_ (default)	[[cpu]] (default)	strict(cpu) (default)	隐含在主机编译中
主机+设备功能	__host__ __device__	__host_ __device__	[[hc]] [[cpu]]	restrict(amp,cpu)	没有等价
内核启动	<<< >>>	hipLaunchKernel	hc:: parallel_for_each	concurrency:: parallel_for_each	clEnqueueND- RangeKernel
全局内存	__global__	__global__	不必要/隐含	不必要/隐含	__global
组内存	__shared__	__shared__	tile_static	tile_static	__local
常数	__constant__	__constant__	不必要/隐含	不必要/隐含	__constant
	__syncthreads	__syncthreads	tile_static.barrier()	t_idx.barrier()	barrier(CLK_LOCAL_MEMFENCE)
原子内置	atomicAdd	atomicAdd	hc::atomic_fetch_add	concurrency:: atomic_fetch_add	atomic_add
精确数学	cos(f)	cos(f)	hc:: precise_math::cos(f)	concurrency:: precise_math::cos(f)	cos(f)
快速数学	__cos(f)	__cos(f)	hc::fast_math::cos(f)	concurrency:: fast_math::cos(f)	native_cos(f)
向量	float4	float4	hc:: short_vector::float4	concurrency:: graphics::float_4	float4

对于HC和C++AMP，假设捕获的_tiled_ext_名为t_ext，捕获的_extent_称为ext。这些语言使用捕获的变量将信息传递给内核，而不是使用特殊的内置函数，因此确切的变量名称可能会有所不同。

索引函数（从线程索引开始）显示了1D网格的术语。一些API对3D网格使用xyz/012索引的逆序。

HC允许在运行时指定图块维度，而C++AMP要求在编译时指定图条维度。因此，拼贴块dims的hc语法是t_ext.tile_dim[0]，而C++AMP是t_ext.teile_dim0。

从ROCm 2.0版本开始，C++AMP在HCC中不再可用。