GPU编程和流式多处理器

流式多处理器（SM）是运行CUDA内核的GPU的一部分。本章重点介绍SM的指令集功能。

流式多处理器（SM）是运行我们的CUDA内核的GPU的一部分。每个SM包含以下内容。

• 可以在执行线程之间划分的数千个寄存器
• 几个缓存：
  • –共享内存，用于线程之间的快速数据交换
  • –恒定高速缓存，用于快速广播恒定内存中的读取
  • –纹理缓存，以聚合纹理内存的带宽
  • – L1_缓存，_可减少对本地或全局内存的延迟
• Warp_调度程序可以快速在线程之间切换上下文，并向准备执行的_Warp发出指令
• 用于整数和浮点运算的执行核心：
  • –整数和单精度浮点运算
  • –双精度浮点
  • –用于单精度浮点先验功能的特殊功能单元（SFU）

存在许多寄存器，以及硬件可以如此高效地在线程之间进行上下文切换的原因，可以最大程度地提高硬件的吞吐量。GPU被设计为具有足够的状态，以覆盖执行等待时间和数百个时钟周期的存储等待时间，执行读取指令后，设备存储器中的数据可能需要数百个时钟周期才能到达。

SM是通用处理器，设计与CPU中的执行内核有很大不同：目标时钟频率低得多；支持指令级并行性，但不支持分支预测或推测性执行；如果根本没有缓存，则缓存较少。对于合适的工作负载，GPU中强大的计算能力足以弥补这些缺点。

自2006年推出首款支持CUDA的硬件以来，SM的设计一直在迅速发展，其代号为Tesla，Fermi和Kepler的三个主要修订版本。开发人员可以通过调用cudaGetDeviceProperties（）并检查cudaDeviceProp.major和cudaDeviceProp.minor或通过调用驱动程序API函数cuDeviceComputeCapability（）来查询计算能力。计算能力1.x，2.x和3.x分别对应于Tesla类，Fermi类和Kepler类硬件。表8.1总结了每一代SM硬件中添加的功能。

表1 SM功能

以后会显示不同SM的框图。CUDA内核可以执行整数和单精度浮点指令。一个双精度单元实现双精度支持（如果可用）；和特殊功能单元实现倒数，倒数平方根，正弦/余弦和对数/指数函数。当执行指令所需的资源可用时，warp调度程序会将指令调度到这些执行单元。

本文重点介绍SM的指令集功能。因此，有时会引用“ SASS”指令，即ptxas或CUDA驱动程序将中间PTX代码转换成的本机指令。开发人员无法直接编写SASS代码；相反，NVIDIA通过cuobjdump实用程序使这些说明对开发人员可见，可以通过检查编译的微代码来指导其源代码的优化。

1.1. 寄存器

每个SM包含数千个32位寄存器，这些内核在启动内核时指定给线程分配。寄存器是SM中最快，最丰富的内存。例如，开普勒类（SM 3.0）SMX包含65,536个寄存器或256K，而纹理缓存仅为48K。

CUDA寄存器可以包含整数或浮点数据。对于能够执行双精度算术（SM 1.3和更高版本）的硬件，操作数包含在偶数值寄存器对中。在SM 2.0和更高版本的硬件上，寄存器对也可以保存64位地址。

CUDA硬件还支持更广泛的内存事务：内置int2 / float2和int4 / float4数据类型分别位于对齐的寄存器对或四元组中，可使用单个64位或128位宽的加载或存储来读取或写入。一旦进入寄存器，就可以将各个数据元素称为.x / .y（对于int2 / float2）或.x / .y / .z / .w（对于int4 / float4）。

通过指定命令行选项--ptxas-options --verbose，开发人员可以使nvcc报告内核使用的寄存器数。内核使用的寄存器数量会影响SM中可容纳的线程数量，通常必须仔细调整，以获得最佳性能。可以使用--ptxas-options --maxregcount N指定用于编译的最大寄存器数。

注册采样Register Aliasing

寄存器可以保存浮点或整数数据，某些内在函数仅用于迫使编译器更改其变量视图。__int_as_float（）和__float_as_int（）内部函数，导致32位整数和单精度浮点之间“改变性能”的变量。

float__int_as_float（int i）;

int __float_as_int（float f）;

__double2loint（） ，__double2hiint（） ，和__hiloint2double（）内部函数类似原因寄存器来改变个性（通常就地）。__double_as_longlong（）和__longlong_as_double（）就地强制寄存器对；__double2loint（）和__double2hiint（）分别返回输入操作数的最低和最高32位；__hiloint2double（）从高半部分和低半部分中构造一个双精度型。

int double2loint（double d）;

int double2hiint（double d）;

int hiloint2double（int hi，int lo）;

double long_as_double（long long int i）;

long long int __double_as_longlong（double d）;

1.2. 本地Local Memory

本地存储器用于溢出寄存器，还用于存储已索引且无法在编译时计算其索引的局部变量。本地内存由与全局内存相同的设备内存池支持，因此它具有与Fermi和更高版本的硬件上的L1和L2缓存层次结构相同的延迟特性和优势。本地内存的寻址方式是自动合并内存事务。硬件包括加载和存储本地内存的特殊说明：SASS变体是Tesla的LLD / LST和Fermi和Kepler的LDL / STL。

1.3. 全局Global Memory

SM可以使用GLD / GST指令（在Tesla上）和LD / ST指令（在Fermi和Kepler上）读取或写入全局内存。开发人员可以使用标准的C运算符来计算和取消引用地址，包括指针算法和取消引用运算符*，[]和->。对64位或128位内置数据类型（int2 / float2 / int4 / float4）进行操作，自动使编译器发出64位或128位加载和存储指令。通过合并内存事务可实现最大的内存性能。

特斯拉级硬件（SM 1.x）使用​​特殊的地址寄存器来保存指针。后来的硬件实现了一种加载/存储架构，该架构使用相同的寄存器文件来存储指针。整数和浮点值；以及用于恒定内存，共享内存和全局内存的相同地址空间。1个

费米级硬件包括旧硬件不具备的一些功能。

• 通过“宽”加载/存储指令支持64位寻址，其中地址保存在偶数寄存器对中。在32位主机平台上不支持64位寻址。在64位主机平台上，将自动启用64位寻址。结果，针对为32位和64位主机平台编译的，相同内核生成的代码，可能具有不同的寄存器计数和性能。
• L1缓存的大小可以配置为16K或48K。2（Kepler添加了将缓存拆分为32K L1 / 32K共享功能。）加载指令可以包括可缓存性提示（告诉硬件将读取的内容拖入L1或绕过L1，并将数据仅保留在L2中）。可通过嵌入式PTX或通过命令行选项–X ptxas –dlcm = ca（默认设置在L1和L2中缓存）或–X ptxas –dlcm = cg（仅在L2中缓存）访问这些文件。

即使多个GPU线程在同一内存位置上运行，原子操作（或仅仅是“原子”）也可以正常工作地更新内存位置。在操作期间，硬件会在内存位置强制执行互斥。由于不能保证操作顺序，因此通常支持的运算符是关联的。3

Atomics首先可用于SM 1.1和更高版本的全局内存以及SM 1.2和更高版本的共享内存。但是，在开普勒一代硬件出现之前，全局内存原子太慢而无法使用。

当通过--gpu-architecture为nvcc指定了适当的体系结构时，表2中汇总的全局原子内在函数将自动变为可用。所有这些内在函数都可以对32位整数进行操作。SM 1.2中添加了对atomicAdd（），atomicExch（）和atomicCAS（）的64位支持。在SM 2.0中添加了32位浮点值（float）的atomicAdd（）。在SM 3.5中添加了对atomicMin（），atomicMax（），atomicAnd（），atomicOr（）和atomicXor（）的64位支持。

表2原子操作

注意

由于原子操作是使用GPU的集成内存控制器中的硬件实现的，无法在PCI Express总线上运行，无法在与主机内存或对等内存相对应的设备内存指针上正确运行。

在硬件级别，原子有两种形式：原子操作返回执行算子之前在指定内存位置的值，以及归约操作，开发人员可以在内存位置“触发并忘记”而忽略返回值值。由于如果不需要返回旧值，则硬件可以更有效地执行操作，编译器将检测是否使用了返回值，如果未使用，则发出不同的指令。例如，在SM 2.0中，指令分别称为ATOM和RED。

1.4. 恒定内存Memory

常量内存驻留在设备内存中，但由另一个只读缓存支持，该缓存经过优化，可以将读取请求的结果广播，到均引用同一内存位置的线程。每个SM均包含一个经过延迟优化的小型缓存，用于处理这些读取请求。将内存（和缓存）设置为只读可简化缓存管理，因为硬件无需实施回写策略来处理已更新的内存。

SM 2.x和后续硬件包括针对内存的特殊优化，该优化未表示为常量，但编译器已将其标识为（1）只读和（2）其地址不依赖于块或线程ID。“均匀加载”（LDU）指令使用恒定的缓存层次结构读取内存，并将数据广播到线程。

1.5. 共享内存

共享内存非常快，是SM中的片上内存，线程可以将其用于线程块内的数据交换。由于它是每个SM资源，因此共享内存的使用会影响占用率，即SM可以保留的warp数量。SM使用特殊指令加载和存储共享内存：SM 1.x上为G2R / R2G，SM 2.x及更高版本上为LDS / STS。

共享内存安排为交错的存储体，通常针对32位访问进行了优化。如果经线中有多个线程引用同一存储体，则会发生存储体冲突，并且硬件必须连续处理内存请求，直到为所有请求提供服务为止。为了避免存储区冲突，应用程序将基于线程ID以交错模式访问共享内存，如下所示。

extern __shared__ float shared[];

float data = shared[BaseIndex + threadIdx.x];

从相同的32位共享内存位置读取扭曲中的所有线程也很快。硬件包括广播机制以针对这种情况进行优化。写入同一存储区的操作会由硬件进行序列化，降低性能。写入同一地址会导致争用情况，应避免。

对于2D访问模式（例如图像处理内核中的像素图块），最好填充共享内存分配，以便内核可以引用相邻行，而不会引起存储体冲突。SM 2.x和后续硬件具有32个存储区，其中每个存储区可供2个图块使用，其中同一warp中的线程可以按行访问数据，这是将图块大小填充为33个32位字的倍数的好策略。

在SM 1.x硬件上，共享内存的大小约为16K。在更高版本的硬件上，总共有64K的L1高速缓存，可以配置为16K或48K共享内存，其余部分用作L1高速缓存。

在过去的几代硬件中，NVIDIA改进了硬件对除32位以外的操作数大小的处理。在SM 1.x硬件上，来自同一存储区的8位和16位读取导致存储区冲突，而SM 2.x和更高版本的硬件可以从同一存储区广播中，任何大小的读取。同样，共享内存中的64位操作数（例如double）比SM 1.x上的32位操作数慢得多，以至于开发人员有时不得不求助于将数据分别存储为上下半部分。SM 3.x硬件为主要在共享内存中使用64位操作数的内核添加了一项新功能：将存储体大小增加到64位的模式。

共享内存中的原子算子

SM 1.2添加了在共享内存中执行原子操作的功能。全局内存与使用单个指令（取决于GATOM或GRED，取决于是否使用返回值）来实现原子的全局存储器不同，共享内存原子是通过显式的锁定/解锁语义实现的，编译器发出的代码会导致每个线程循环这些锁定操作，直到线程执行其原子操作为止。

清单1将源代码提供给atomic32Shared.cu，该程序专门用于编译以突出显示共享内存原子的代码生成。清单2显示了为SM 2.0生成的最终微代码。注意，LDSLK（与锁共享负载）指令如何返回谓词，该谓词说明是否获取了锁，确定了要执行更新的代码，并且代码循环运行，直到获取了锁并执行了更新。

锁是按32位字执行的，锁的索引由共享内存地址的位2–9确定。注意避免争用，否则清单2中的循环最多可以迭代32次。

list1. atomic32Shared.cu

__global__ void

Return32( int *sum, int *out, const int *pIn )

{

extern __shared__ int s[];

s[threadIdx.x] = pIn[threadIdx.x];

__syncthreads();

(void) atomicAdd( &s[threadIdx.x], *pIn );

__syncthreads();

out[threadIdx.x] = s[threadIdx.x];

}

Listing 2. atomic32Shared.cubin (microcode compiled for SM 2.0)

code for sm_20

Function : _Z8Return32PiS_PKi

/*0000*/     MOV R1, c [0x1] [0x100];

/*0008*/     S2R R0, SR_Tid_X;

/*0010*/     SHL R3, R0, 0x2;

/*0018*/     MOV R0, c [0x0] [0x28];

/*0020*/     IADD R2, R3, c [0x0] [0x28];

/*0028*/     IMAD.U32.U32 RZ, R0, R1, RZ;

/*0030*/     LD R2, [R2];

/*0038*/     STS [R3], R2;

/*0040*/     SSY 0x80;

/*0048*/     BAR.RED.POPC RZ, RZ;

/*0050*/     LD R0, [R0];

/*0058*/     LDSLK P0, R2, [R3];

/*0060*/     @P0 IADD R2, R2, R0;

/*0068*/     @P0 STSUL [R3], R2;

/*0070*/     @!P0 BRA 0x58;

/*0078*/     NOP.S CC.T;

/*0080*/     BAR.RED.POPC RZ, RZ;

/*0088*/     LDS R0, [R3];

/*0090*/     IADD R2, R3, c [0x0] [0x24];

/*0098*/     ST [R2], R0;

/*00a0*/     EXIT;

……………………………..

1.6. 障碍和连贯性

熟悉的__syncthreads（）内部函数会等待，直到继续执行线程块中的所有线程为止。需要保持线程块内共享内存的一致性。其他类似的内存屏障指令，也可以用于在更大范围的内存上执行某些排序，如表3所述。

表3内存屏障本质