[C++知识库]CUDA C 编程权威指南 Grossman 第7章 调整指令级原语

为了在GPU上实现较高吞吐量，你需要了解有哪些因素限制峰值的性能：延迟，带宽或算术运算。基于此，可以将应用程序分为两类：

????????I/O密集型；

????????计算密集型；

处理器的计算吞吐量可以用它在一段时间内执行操作的数量来衡量。因为GPU有很多SIMT指令和计算核心，所以其峰值计算吞吐量通常比其他的处理器要高。

对应用程序的吞吐量和正确性进行优化时，理解不同低级原语的性能，数值精度和线程安全性方面的优缺点是很重要的。知道内核代码在什么时候被编译成一条原语或其他语句，能让你根据需求调整编译器生成的代码。

执行MAD(乘加指令)的结果是循环次数减少了一半。这种性能的提升并不是没有代价。一个MAD指令的数据精确性往往比单独的乘法和加法指令的要低。

在本章中，将学习各种用于优化性能、提高准确性和正确性比较低级的CUDA的原语。

7.1 CUDA指令概述

指令是处理器的一个逻辑单元。知道CUDA内核代码什么时候会产生不同指令以及高级语言如何转换为指令，却是很重要的。对于两个功能等效指令可以影响很多应用程序的特性，包括性能、精确度和正确性。当通过严格的数字请求时，把遗留应用程序传输到CUDA时，就要特别留意这些问题。

显著影响CUDA内核生成指令的3大因素：浮点运算、内置和标准函数、原子操作。

浮点运算是针对非整数值的运算，并且会影响CUDA程序的精确度和性能。

7.1.1 浮点指令

自从浮点运算采用IEEE-754标准后，规定将二进制浮点数编码成3段：符号段s(sign)，一个比特位；指数段e(exponent)，多个比特位；以及尾数v或者分数段(fraction)，多个比特位。

?

?浮点型数值不能精确存储，只能在四舍五入后在存储。

浮点编程需要考虑的另一个方面是浮点数的粒度问题。浮点数粒度比整数来说要好。然而浮点数只能在离散的区间间隔内存储数据。随着浮点数值距离0越来越远，表示数值的区间会随之增大。

可以使用C语言中的数学函数nextafterf，从一个给定值找到下一个最高位可表示数字的区别。值得注意的是，随着x值的增大，精度会大幅降低。

在浮点数值上进行操作的指令被称为浮点指令。

7.1.2 内部函数和标准函数

除了单精度和双精度操作的区别，CUDA还将所有算术函数分成内部函数和标准函数。标准函数用于支持可对主机和设备进行访问并标准化主机和设备的操作。

标准函数包括来自于C标准数学库的数学运算。

CUDA内置函数只能对设备代码进行访问。在编程中，如果一个函数是内部函数或是内置函数，那么在编译时对它的行为会有特殊响应，从而产生更积极的优化和更专业化的指令生成。

在CUDA中，许多内部函数与标准函数是有关联的，这意味着存在与内部函数功能相同的标准函数。举个例子，标准函数中的双精度浮点平方根函数也就是sqrt。有相同功能的内部函数是_dsqrt_rn。还有执行单精度浮点除法运算的内部函数：__fdividef。

内部函数分解成了比与它们等价的标准函数更少的指令。这会导致内部函数比等价的标准函数更快，但数值精确度却更低。因此可以在同一应用中交替使用标准函数和内部函数，但是它们在性能和数值精确度会有所不同。

标准函数和内部函数大大增加了CUDA应用程序的灵活性。它们作为细粒度旋钮，可以在运行操作基础上调整性能和数值精确度。

7.1.3 原子操作指令

一条原子操作指令用来执行一个数学运算，此操作时一个独立不间断的操作，且没有其他线程的干扰。

原子操作指令阻止了多个线程之个互相干扰，他们可以对跨线程共享数据进行“读-改-写”操作。

CUDA提供了在32位或64位全局内存或共享内存上执行读-改-写操作的原子函数。

所有计算能力为1.1及以上设备都支持原子操作。

如果不止一个线程对同一个内存位置进行写操作，这叫做数据竞争。

不止一个线程对同一个内存位置进行写操作，这叫做数据竞争，或者称为对内存的不安全访问。数据竞争的定义时两个或多个独立的正在执行的线程访问同一地址，并且至少其中一个访问会修改该地址。幸好，使用原子操作指令可以避免这种事情的发生。

比如atomicAdd(int* M, int V)，大多数原子函数是二进制函数，能够在两个操作数上进行操作。它们把一个内存位置M和一个数值V作为输入。与原子函数相关的操作在V上执行，数值V早已存储在内存地址M中，然后把运算结果写到同样的内存位置了。

原子运算函数分为3种：算术运算函数、按位运算函数和替换函数。

原子算术在目标内存位置上执行简单的算术运算，包括加、减、最大、最小、自增和自减等操作。

原子按位运算函数在目标内存位置进行按位与，或，按位异或。

原子替代函数可以用一个新值来替换内存位置上原有的值，它可以是有条件的也可以是无条件的。不管替换是否成功，原子替换函数总是会返回最初存储在目标位置上的值。atomicExch可以无条件地替换已有的值。如果当前存储的值与GPU线程调用指定的值相同，那么atomicCAS可以有条件地替换已有的值。

// 对flag的不安全访问
__global__ void check_threadhold(int* arr, int threadhold, int* flag)
{
    if (arr[blockIdx.x* blockDim.x + threadIdx.x] > threadhold)
        *flag = 1;
}

// 用原子操作
__global__ void check_threadhold(int* arr, int threadhold, int* flag)
{
    if (arr[blockIdx.x* blockDim.x + threadIdx.x] > threadhold)
        atomicExch(flag, 1);
}

事实上，使用atomicExch等原子操作可能会显著降低其性能。当使用这种优化时必须要非常小心，因为这种运算并不依赖与每个线程可见的运算结果。

原子操作指令在高并行环境如GPU是很强大的。它们提供了一种安全的方法，来操作被成百上千线程所共享的数据。虽然原子函数没有精确度上的顾虑(而内部函数需要考虑精确度)，但是它们的使用会严重降低性能。

7.2 程序优化指令

用于优化程序的指令，有很多的选择，单精度或双精度浮点数值、标准或内部函数、原子函数或不安全访问。

7.2.1 单精度和双精度比较

用于存储单精度和双精度的位数是不同的。双精度的数值精确性是以空间和性能消耗为代价的。

单精度和双精度浮点运算在通信和计算上的性能差异是不可忽略的。在设备端进行数据通信的时间也是使用单精度数值的两倍，这是由双精度数值长度是单精度数值长度的两倍造成的。随着全局内存输入/输出数量和每条指令执行的位操作数量的增加，设备上的计算时间也会增加。

考虑到数值精确度，在迭代应用中可能更需要使用双精度变量。

在声明单精度浮点数时必须非常谨慎。任何不正确的省略尾数f的声明都会自动地被NVCC编译器转换为双精度数。

小结;

? ? ? ? 浮点运算对应用程序的性能和数值精确度上的影响并不只是在GPU上会产生，使用其他架构时，也会遇见同样的问题。以下是GPU和CUDA独有的特点:

? ? ? ? 使用双精度数值会增加主机和设备之间的通信；

? ? ? ? 使用双精度会增加全局内存的输入/输出；

? ? ? ? 数值精度的损失是由CUDA编译器强制浮点数值优化导致的；

7.2.2 标准函数与内部函数的比较

7.2.2.1 标准函数和内部函数可视化

使用nvcc的--ptx标志能够让编译器在并行线程执行(PTX)和指令集架构(ISA)中生成程序的中间表达式，而不是生成一个一个最终的可执行文件。PTX类似于x86编程里面的程序集，它提供了一个你所编写的内核代码之间的中间表达式，该指令在GPU上执行。它对于深入了解内核的低级别执行路径是很有用的。

可以用文本编辑器打开.ptx文件。

.entry指令标志了一个函数定义的开始。对于pow函数而言，标准函数签名__Z8standardPf，对于内置函数而言，__Z8intrinsicPf。函数签名可能会因编译器版本的不同而不同。对于powf函数而言，内部函数实现只需17行，而标准函数(cuda5.0)使用了344行。

区分标准函数和内部函数的不仅有性能，它们的计算精度也是不同的。使用内置函数，性能可能提升24倍。精度差了一个数量级。

CPU到GPU的移植：

? ? ? ? 需要说明允许的误差范围；

7.2.2.2 操纵指令生成

在大多数情况下，将程序员编写的内核代码转换为GPU指令集这一过程是由CUDA编译器完成的。程序员很少会有检查或手动修改指令的想法。但是，你也可以引导编译器倾向于实现良好的性能或准确性或者达到二者的平衡。

CUDA编译器有两种方法可以控制指令级优化类型：编译器标志、内部或标准函数调用。

一个个手动调整内核操作的工作量太大了。编译器标志提供了一个更自动、全局化的方式来操从编译器指令的生成。nvcc的--fmad选项可全局地启用或禁用FMAD整个编译单元的优化。

除了--fmad，还有许多CUDA编译器标志会影响算法指令的生成。完整的列表可在nvcc --help选项中找到。

在有*乘法运算法的地方调用__fmul或dmul时，将阻止nvcc将乘法作为乘加优化的一部分来使用。因此可以通过有选择性地调用__fmul或者__dmul的计算来提高某些数值的健壮性时，可启用MAD编译器优化全局。

许多浮点型内部函数包括(__fadd, __fsub, __fmul等)在函数名中都使用两个后缀字符，者明确指出浮点四舍五入的模式。

7.2.2.3 小结

7.2.3 了解原子指令

在本节中，你将学习如何使用原子操作，并学习在高并发环境下的共享数据如何执行正确的操作。

7.2.3.1 从头开始

通过使用一个原子函数，每个由CUDA提供的原子函数可以重复被执行：原子级比较并交换CAS运算符。原子级CAS是一个很重要的操作，不仅可以使你在CUDA中定义你自己的原子函数，还能帮助你更深层次地理解原子操作。

CAS将3个内容作为输入：内存地址，存储在此地址中的期望值，已经实际想要存储在此位置的新值；然后执行以下几步;

? ? ? ? 1. 读取目标地址并将该处地址的存储值与预期值进行比较。

? ? ? ? ? ? ? ? a. 如果存储值与预期值相等，那么新值将存入目标位置。

? ? ? ? ? ? ? ? b. 如果存储值与预期值不等，那么目标位置不会发生变化。

? ? ? ? 2. 不论发生声情况，一个CAS操作总是返回目标地址中的值。注意，使用返回值可以来检查一个数值是否被替换成功。如果返回值等于传入的预期值，那么CAS操作一定成功了。

atomic(int* address, int compare, int val);address是目标内存地址，compare是预期值，val是实际想写入的值。

如何利用atomicCAS执行一个原子加法？ 首先需要分解加法运算并把它定义成CAS操作。当执行自定义原子操作时，定义目标的起始和结束状态是很有帮助的。在原子加法中，起始状态是递增运算符的基值。结束状态是起始状态和增量的总和。这个定义直接转换为atomicCAS：预期值是起始状态，实际写入的是新值是完成状态。

// 此版本的问题是：目标地址是由多线程共享的，所以另一个线程修改address是可能的，这个值
// 处于被expected读入和atomicCAS修改之间。
__device__ int myAtomicAdd(int* address, int incr)
{
    int expected = *address;
    int oldValue = atomicCAS(address, expected, expected + incr);
}

回忆一下，如果atomicCAS的返回值与预期值不同则程序会失败。因此，myAtomicAdd可以用来检查失败并在一个循环中重试CAS直到atomicCAS成功。

__device__ int myAtomicAdd(int* address, int incr)
{
    int expected = *address;
    int oldValue = atomicCAS(address, expected, expected + incr);

    while (oldValue != expected)
    {
       expected = *address;
       oldValue = atomicCAS(address, expected, expected + incr);
    }

    return oldValue;
}

7.2.3.2 内置的CUDA原子函数

对共享内存32位操作和全局内存64位操作数值支持始于计算能力1.2的设备。对共享内存64位数值的操作支持始于计算能力2.0设备。仅有atomicAdd支持float和double类型。

原子函数

7.2.3.3 原子操作成本

原子函数导致很高的性能代价的原因如下：

? ? ? ? 1. 当在全局或共享内存中执行原子操作时，能保证所有的数值变化对所有线程都是立即可见的。因此，在最低限度下，一个原子操作指令将通过任何方式进入到全局或共享内存中读取当前存储的数值而不需要缓存。如果原子操作指令成功，那么必须吧把实际需要的值写入到全局或共享内存。

? ? ? ? 2. 共享地址冲突的原子访问可能要求发生冲突的线程不断地进行重试，类似于运行多个myAtomicAdd循环的迭代。如果你的应用程序反复循环而致使IO开销较大，相应地性能会降低。

? ? ? ? 3. 当线程在同一个线程束中时必须执行不同的指令，线程束的执行是序列化的。如果一个线程束中的多个线程在相同的内存地址发出一个原子操作，那么产生类似于线程冲突的问题。因为只有一个线程的原子操作可以成功，所有其他的线程必须重试。并且一个原子操作也意味着一次全局的读取和写入。

7.2.3.4 限制原子操作的性能成本

有些方法能够减少原子操作带来的性能损失：

? ? ? ? 可以使用局部操作来增强全局原子操作，这些局部操作能从同一个线程块的线程中产生一个中间结果。哲学要使用本地较低延迟的资源，如shuffle指令或共享内存，在使用原子操作把局部结果结合到最终的全局结果之前，需要先重每个线程块产生局部结果。

7.2.3.5 原子级浮点支持

原子函数中要注意一点是它们大多被声明在整型数值上操作。只有atomicExch和atomincAdd支持浮点数值。

在高级别上，有一个办法是用一个变量中支持的类型存储浮点数的原始比特位，并使用所支持的类型执行原子CAS操作。

下面是一个用单精度浮点数实现myAtomicAdd核函数;

__device__ float myAtomicAdd(float* address, float incr)
{
    unsigned int* typeAddress = (unsigned int*)address;

    float currentVal = *address;

    unsigned int expected = __float2uint_rn(currentVal);

    unsigned int desired = __float2uint_rn(currentVale + incr);

    int oldIntValue = atomicCAS(typeAddress, expected, desired);

    while(oldIntValue != expected)
    {
        expected = oldIntValue;

        desired = __float2uint_rn(__uint2float_rn(oldIntValue) + incr);
        oldIntValue = atomicCAS(typeAddress, expected, desired);

    }

    return __uint2float_rn(oldIntValue);
}

此核函数使用了：

? ? ? ? 1. 一个cast改变了address指针的类型，使其从float到unsigned int型；

? ? ? ? 2. 使用__float2uint_rn将期望值、*address、期望值以及*address + incr的类型转换为包含相同比特位的unsigned int类型；

? ? ? ? 3. 如果操作失败了，使用__uint2float_rn检索一个从atomicCAS返回的unsigned int浮点数并计算新的期望值。

? ? ? ? 所有的这些类型转换是必要的，因为应用程序要求的类型(float)和atomicCAS函数要求的类型是不同的。CUDA提供了一个有很大范围的有其他特定累心转换的函数，包括__double_as_longlong、__longlong_as_double、__douoble-2float_rn等。详细的可以在CUDA Mah API文档中找到。

7.2.3.6 小结

使用双精度导致性能下降的原因主要是：

? ? ? ? 1.由于双精度的长度是单精度类型的两倍，所以导致两倍主机-设备通信开销；

? ? ? ? 2. 由于从全局内存加载了两次数据，增加了设备上IO开销；

? ? ? ? 3. 双精度导致能存入的寄存器的更少，同时每个线程块中每个线程可用资源的减少，潜在地导致全局内存中有更多溢出变量；

? ? ? ? 4. 在两倍比特位上执行算术运算增加了计算成本。

7.3 总结

? ? ? ? 在写CUDA时，需要明白一下几点性能和数值精度的影响：浮点型运算，标准和内部函数以及原子操作。

? ? ? ? CUDA编译器和函数库通常隐藏了底层细节，这是一把双刃剑。自动编译器的优化减少了一些优化负担，但可能会导致内核中数据转化变得不可见。这种不透明性导致数值问题调试困难。