第 11 章：关于标量化和亲和执行的 GPU 研究方向

如第 2 章所述，GPU 计算 API（如 CUDA 和 OpenCL）采用了类似 MIMD 的编程模型，允许程序员在 GPU 上启动大量标量线程。虽然这些标量线程中的每一个都可以遵循其独特的执行路径并访问任意内存位置，但在常见情况下，它们都遵循少量执行路径并执行相似的操作。

GPU 线程之间的收敛控制流在大多数（如果不是全部）现代 GPU 上都通过 SIMT 执行模型进行利用，其中标量线程被分组为在 SIMD 硬件上运行的线程束（参见第 3.1.1 节）。本章总结了一系列进一步利用这些标量线程相似性的研究，包括标量化和亲和执行。

这些研究的关键洞见在于观察到计算内核中执行的线程之间的值结构[Kim et al., 2013]。均匀和亲和两种值结构类型在示例 11.1 中的计算内核中进行了说明。

均匀变量 对于计算内核中的每个线程，都具有相同的常量值的变量。在算法 11.1 中，变量 a 以及字面量 THRESHOLD 和 Y_MAX_VALUE 都具有在计算内核中的所有线程中均匀的值。均匀变量可以存储在单个标量寄存器中，并被计算内核中的所有线程重复使用。

亲和变量 对于计算内核中的每个线程，其值是线程 ID 的线性函数的变量。在算法 11.1 中，变量 y[idx] 的内存地址可以表示为线程 ID threadIdx.x 的亲和变换：

&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;

这种亲和表示可以存储为一对标量值（基址和步长），这比完全展开的向量更加紧凑。

__global__ void vsadd( int y[], int a ) {
    // 在此处添加代码
}
```以下是该 Markdown 文件的中文翻译版本。对于代码部分,只翻译注释,代码本身不翻译。
 
```c
int idx = threadIdx.x;
    y[idx] = y[idx] + a;
    if ( y[idx] > THRESHOLD )
        y[idx] = Y_MAX_VALUE;
}

算法 11.1: 计算内核中标量和仿射操作的示例 (来自 [Kim et al., 2013])。

有多项研究提出如何在 GPU 上检测和利用统一或仿射变量。本章总结了这两个方面的这些提议。

检测统一或仿射变量

检测 GPU 计算内核中统一或仿射变量的存在有两种主要方法:编译器驱动检测和硬件检测。

编译器驱动检测

检测 GPU 计算内核中统一或仿射变量存在的一种方法是通过特殊的编译器分析。这是可能的,因为现有的 GPU 编程模型 CUDA 和 OpenCL 已经为程序员提供了将变量声明为整个计算内核中常量的方法,以及提供了一个特殊的线程 ID 变量。编译器可以执行控制依赖分析,检测那些仅依赖于常量和线程 ID 的变量,并将它们标记为统一/仿射。仅在统一/仿射变量上工作的操作然后成为标量化的候选对象。

AMD GCN [AMD, 2012] 依赖于编译器来检测统一变量和可以由专用标量处理器存储和处理的标量操作。

Asanovic 等人 [2013] 引入了一种结合收敛和变体分析的方法,允许编译器确定任意计算内核中合格进行标量化和/或仿射变换的操作。计算内核的收敛区域内的指令可以转换为标量/仿射指令。在从发散区域到收敛区域的任何转换点,编译器都会插入一个 syncwarp 指令来处理两个区域之间由控制流引起的寄存器依赖关系。Asanovic 等人 [2013] 采用了这种方法。以下是该 Markdown 文件的中文翻译。对于代码部分,请不要翻译代码,只翻译注释。

这个分析用于为 Temporal-SIMT 架构 [Keckler 等人, 2011, Krashinsky, 2011] 生成标量操作。

分离式仿射计算 (DAC) [Wang 和 Lin, 2017] 依赖于类似的编译器分析来提取标量和仿射候选项,将其分离到单独的 warp 中。Wang 和 Lin [2017] 增加了发散仿射分析的过程,目的是提取指令序列。