SM 处理一条 warp 指令，首先为 warp 里的每个线程读指令操作数，执行指令，最后为 bank
half-warp
bank
half-warp
warp 里的每个线程写入计算结果。因此，有效的指令吞吐量不仅取决于名义上的指令吞吐量，
还取决于内存延迟和带宽。建议采取以下手段增大指令吞吐量：
避免使用低吞吐量指令。
对每种类型的存储器进行优化，有效利用带宽。
允许线程调度单元尽量用多的数学计算来覆盖访存延迟，这就需要有高的算术密度
（或者说，对于每一个内存操作都有大量的算术操作来覆盖），同时每个多处理器有 很多的活动线程。
本节中，吞吐量是指每个多处理器在一个时钟周期下执行的操作数目。对于大小为 32 的 warp，一条指令由 32 个操作构成。因此，如果记 T 为每个时钟下的操作数目，那么指令吞吐 量就是每 32/T 个时钟周期一条指令。
所有的吞吐量都是针对一个多处理器而言的。所以，要计算整个设备的吞吐量需要乘以 设备的多处理器个数。
对程序中需要多次运行的代码进行指令级优化可以有效地提高指令流的吞吐量。本节将 对各种指令进行详细讨论。
4.5.1 算术指令
CUDA 的硬件特别适合进行单精度浮点运算，因此应该尽量使用单精度浮点单元进行计算。
如果在代码中可以使用单精度浮点代替双精度浮点，那么我们强烈建议使用 float 型和单 精度浮点数学函数。如果代码中使用了双精度，那么将代码编译到不支持双精度的硬件设备上 时，例如计算能力为 1.2 或更低的设备，每个双精度的变量将会转成单精度格式（但大小仍是 64bit），并且双精度算术运算也会转为单精度算术运算。
下面是算术运算的吞吐量和使用时需要注意的问题，以及优化方法。
4.5.1.1 单精浮点基本算术运算
单精浮点加、乘、乘加运算的吞吐量是每个时钟周期 8 个操作。
求倒数运算的吞吐量是每个时钟周期 2 个操作。
单精浮点除操作是每个时钟周期 0.88 个操作，但是__fdividef(x, y)提供一个更快速的版本，
能达到每个时钟周期 1.6 个操作。
4.5.1.2 单精浮点的平方根和倒数平方根
倒数平方根的吞吐量是每个时钟周期 2 个操作。
单精浮点平方根的计算方法是求倒数平方根的倒数，而不是在倒数平方根后做乘法，这 是为了在处理 0 和无穷大时能计算得到正确的结果。因此，它的吞吐量是每个时钟周期 1 个 操作。
4.5.1.3 单精浮点的对数运算
__logf(x)的吞吐量是每个时钟周期 2 个操作。
4.5.1.4 正弦和余弦运算
__sinf(x), __cosf(x),__ exp(x)的吞吐量是每个时钟周期 1 个操作。
sinf(x), cosf(x), tanf(x), sincosf(x)和相应的双精指令开销非常昂贵，尤其是 x 绝对值较大时 更是如此，此时就要将 x 的绝对值减小，称为归约操作。
归约代码（参见 math_functions.h 中的实现）由两个代码路径组成，快路径和慢路径。快 路径适用于参数较小的情况，它本质上是一些乘加操作。慢路径适用于参数较大的情况，包含 一系列的计算，以在整个参数范围内求得正确结果。
目前，三角函数中用的是快路径，单精的话参数要小于 48039.0，双精的话要小于 2147483648.0。
慢路径比快路径要求更多的寄存器。为减小寄存器压力，归约操作可能会使用延迟较高 的 local memory，速度会进一步降低。目前，在单精的函数中用了 28bytes 的 local memory，
在双精函数中用了 44bytes，当然这个数字也不一而论。由于慢路径中冗长的计算和 local memory 的使用，所以使用慢路径的三角函数吞吐量会慢一个数量级。
需要同时计算 sin 和 cos 值时，可以使用 sincos 系列函数节约时间，包括：
高速版本的单精度浮点__sincosf()函数。
高精度的单精度浮点 sincosf()函数。
双精度浮点的 sincos()函数。
4.5.1.5 整数算术运算
整数加法的吞吐量是每个时钟周期 8 个操作。
32bit 整数乘的吞吐量是每个时钟周期 2 个操作，但是__mul24 和__umul24 提供有符号和 无符号的 24 位整数乘法，吞吐量可以达到每个时钟周期 8 个操作。在未来的架构中，__[u]mul24 将会比 32it 整数乘更慢，应该为不同版本硬件编译不同的程序。
整数除和模运算开销特别大，应尽量地避免或用位运算代替。例如如果 n 是 2 的幂次方，
那么(i/n)与(i>>log2(n))是等价的，(i%n)与(i&(n-1))是等价的，如果 n 在程序中已经固定的就是 2 的幂次方，那么编译器将会自动进行这些转换。
4.5.1.6 比较运算
比较、min、max 操作的吞吐量是每时钟周期 8 个操作。
4.5.1.7 位运算
任何位运算的吞吐量是每时钟周期 8 个操作。
4.5.1.8 类型转换
类型转换的吞吐量是每时钟周期 8 个操作。
有些时候，编译器会插入一些转换指令，也就引入了一些额外的执行周期。像这种情况：
操作于 char 或 short 上的函数，这些操作数一般需要转换成 int。
双精浮点常量（定义的时候没有任何后缀）作为单精浮点计算的输入。
最 后 这 种 情 况 能 通 过 使 用 单 精 浮 点 常 量 解 决 ， 声 明 变 量 的 时 候 加 上 后 缀 ， 例 如 3.141592653589793f、1.0f、0.5f。
4.5.2 控制流指令
控制流指令（if、switch、do、for、while）可能引起一个 warp 内的线程跳转到不同的分 支，这将严重地影响指令吞吐量。一旦发生分支，那么不同的执行路径就必须被串行地执行，
导致这个 warp 中指令总数增多。当所有分支的指令都执行结束以后，这些线程才会重回到同 一条执行路径上。
在控制流只与线程的 ID 有关时，各 warp 在 block 中的分布是确定的，应该修改控制条件，
尽量避免在 warp 内发生分支。例如，当控制条件只取决于(threadIdx/warp size)时， warp 内就 不会出现分支，因为此时控制条件是严格按照 warp 对齐的。
有的时候，编译器可能会展开循环或通过谓词执行优化 if、switch 语句，这种情况下，warp 也不会分支。程序员同样能够使用#pragma unroll 指令控制 loop unrolling。
谓词执行是编译器与硬件共同完成的：编译器将几个分支中的指令作为一个代码块处理，
并将其中的指令与谓词关联；而硬件在执行时则根据谓词有选择地执行这些语句，从而将控制 依赖转换为数据依赖。使用谓词执行时，依赖于控制条件的指令并不会被跳过。在编译阶段，
这些依赖于控制的指令都会与一个值为 true 或 false 的谓词相关联，在执行时，尽管分支中的 指令仍然会被调度执行，但只有谓词为真的指令才会真正被执行。如果指令在某条分支中的谓 词为假，虽然这条指令还是要跟着预测值为真的指令一起执行，但它不用将结果写回，也不用 计算访存地址或者读取操作数。
在由分支条件控制的指令数目小于等于某个阈值的情况下，编译器会自动使用谓词指令 替换分支指令。如果编译器认为某个条件有可能造成很多存在分支的 warp，那么阈值是 7，否 则阈值是 4。
4.5.3 访存指令
访存指令包括任何读写 shared/local/global memory 的指令。对 local memory 的访问只有在 寄存器不够用或者编译器无法解析地址时才会发生。
由于寄存器文件的大小、shared memory 每 bank 的宽度，以及对显存进行合并访问时达到 最大带宽的访问宽度都是 32bit，因此在读写时将较大的数据（如 float3，double）拆分成每线 程 32bit，或者将多个[u]char 或[u]short 合并成每线程 32bit 的形式访问，可以获得更高的性能。
存储操作的吞吐量是每个时钟周期 8 个操作。当访问 local/global memory 时，会有额外的 400～600 个时钟周期的访问延迟。
举个例子，如下赋值操作的吞吐量：
__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];
对 global memory 进行的读操作的吞吐量是每个时钟周期 8 个操作，在进行 shared memory 读、写的时候也是每个时钟周期 8 个操作，然而读、写 global memory 时还有 400～600 个时 钟周期的访存延迟。
如果有足够多的算术指令可以在访存期间被发射、执行，那么这些执行的时间就可以隐 藏大部分 global memory 访存延迟。不过，总的来说，对显存的访问还是越少越好。
4.5.4 同步指令
在每个线程都不用等待任何其他线程的情况下，__syncthreads()的吞吐量是每时钟周期 8 个操作。