CUDA调整指令级原语-编程知识

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

对应用程序的吞吐量和正确性进行优化时，理解不同低级原语的性能、数值精确度和线程安全性方面的优缺点很重要。

CUDA指令

指令是处理器中的一个逻辑单元。需要了解CUDA内核代码什么时候会产生不同的指令以及高级语言如何转化为指令很重要。

浮点指令

IEEE-754标准定义了32位和64位浮点格式。标准规定将二进制浮点数编码成三段：符号段，1比特；指数段，多比特；以及尾数段，多比特。

然而浮点数的精度是有限的，举个栗子

a，b两个值都不能在float中精确的存储，都只能近似保存，这样两者恰好相等。

在浮点数值上进行操作的指令被称为浮点指令。CUDA支持所有在浮点数值上的常见数学运算。CUDA编程模式也遵守IEEE-754标准，支持两种精度的浮点数值。

内部函数和标准函数

CUDA将所有算数函数分成内部函数和标准函数。标准函数用于支持可对主机和设备进行访问并标准化主机和设备的操作。标准函数包含来自于C标准数学库的数学运算。

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

在CUDA中，许多内部函数与标准函数是有关联的，意味着存在于内部函数功能相同的标准函数。内部函数分解成了比与它们等价的标准函数更少的指令。导致内部函数比等价的标准函数更快，但数值精度更低。

原子操作指令

一条原子指令用来执行一个数学运算，此操作是一个独立不间断的操作，且没有其他线程的干扰。当一个线程在一个变量上成功完成一个原子操作，那么不管有多少线程正在访问这个变量，这个变量的状态都已经发生了改变。在GPU的高并发环境中，保证“读-改-写”操作的完整性非常重要。CUDA提供了在全局内存或共享内存上执行“读-改-写”操作的原子函数。

与标准函数和内部函数类似，每个原子函数可以实现一个基本数学运算。不同于其他类型指令的是，在原子操作指令中，当两个竞争线程共享的内存空间进行操作时，会有一个定义好的行为。

举个栗子

__global__ void incr(int *p){int temp = *p;temp = temp + 1;*p = temp;
}

如果运行这个核函数，在多线程并行环境中，结果是不确定的。不止一个线程对同一个内存位置进行写操作，叫做数据竞争，或者称为对内存的不安全访问。数据竞争是指，多个独立的正在执行的线程访问同一个地址，而且至少有一个访问会修改该地址。

使用原子操作指令可以避免这种情况的发生。原子操作是通过CUDA API访问的函数。例如，

int atomicAdd(int *M, int V);

M是进行原子操作的地址，v是要加上的值。该原子操作将V加到M地址的变量中，并且返回操作之前的值。

另一个函数

int atomicExch(int *m, int v);

无条件的用v替换m中的值，并且返回原先存在m中的值。

程序优化指令

单精度与双精度比较

单精度与双精度浮点运算在通信和计算上的性能差异是不可忽略的。

单精度相较于双精度浮点运算计算和传输更快，但是精度更低。这些结果可能在迭代过程中被不断积累。

标准函数与内部函数比较

使用nvcc的--ptx标志能够让编译器在并行线程执行和指令集架构中生成程序的中间表达式。生成的PTX文件类似汇编的形式，可以直观的了解内核的低级别执行路径。

另外可以用一些编译指令操纵编译器指令的生成。例如，--fmad=false（默认是true）会强制命令编译器禁用混合乘法与加法的优化。具体命令在《CUDA C编程权威指南》表7-3

了解原子指令

通过使用一个原子函数，每个由CUDA提供的原子函数可以重复被执行：原子级比较并交换符（CAS）运算符。

原子级CAS是一个重要的操作，将三个内容作为输入：内存地址、存储在此地址中的期望值，以及实际想要存储在此位置的新值，然后执行

读取目标地址并将该处地址的存储值与预期值比较
如果存储值与预期值相等，那么新值将存入目标位置
如果存储值与预期值不等，那么目标位置不会发生变化
不论发生什么情况，一个CAS操作总是返回目标地址中的值。

一个原子CAS操作意味着整个CAS进程是在没有其他任何线程干扰的情况下完成的。

详细解释一下atomicCAS设备函数

int atomicCAS(int *address, int compare, int val);

进行atomicCAS时，先比较address地址当前的值是否等于compare，如果相等，则把address地址中的值变成val，如果不等，就不变，无论如何都返回比较前address中的值。

通过atomicCAS可以定义自己的原子操作。

原子操作的成本

原子函数在一些应用中很有必要而且很有帮助，但可能要付出很高的性能代价。

主要原因是，如果有多个线程对同一地址进行原子操作，那么会产生类似线程冲突的情况。只有一个线程可以原子操作成功，其他线程必须循环等待。并且，一个原子操作就意味着一个全局的读取和写入。

限制原子操作的成本

可以在使用局部操作来增强全局原子操作。比如从同一个线程块中产生不同的中间结果，在最后进行原子操作写入全局内存。