8.1_内存
8.1 内存
8.1.1 寄存器
每个SM包含了成千上万个32位寄存器，当内核被启动时，这些寄存器会被分配到指定的线程中。在SM中，寄存器是速度最快的，也是数目最多的存储资源。例如，开普勒架构（SM 3.0）的SMX包含65,536个寄存器，容量共256KB，而纹理寄存器只有48KB。
CUDA寄存器可以装入整型或浮点型数据。如果硬件支持的话（SM1.3及更高）可以进行双精度算术运算，操作数被放在偶数寄存器对中。在SM2.0及更高版本的硬件中，一对寄存器同样也可以装入64位地址。
CUDA硬件还支持更宽的内存事务：内置的int2/float2和int4/float4数据类型。它们分别存储在对齐的一对寄存器中或对齐的4个寄存器中，可以通过64或128位宽的单条存取指令读写。只要数据存储在寄存器中，那么数据元素便可以以.x/.y（int2/float2）或.x/.y/.z/.w（int4/float4）形式引用。
开发人员可以使用命令行选项--ptxas和--verbose来让nvcc报告一个内核使用的寄存器数量。一个内核使用的寄存器数量可以影响每
一个SM内可装入的线程数。所以要得到优化的性能，必须仔细的设置。每一次编译使用的最大寄存器数，可以通过--ptexas-options--maxregcount N来指定。
寄存器别名使用
因为寄存器既可以存储整型数据，也可以存储浮点型数据，有些内置函数只能在强制编译器改变对一个变量的解释方式后工作。
__int_as_float()和__float_as_int()这两个内部函数会使变量在32位整型和单精度浮点数之间“改变身份”。
float_int_as_float(int i); int __float_as_int(float f);__double2float()、__double2hanoi()和__hiloint2double()内部函数可以类似地让寄存器数据改变身份（通常原地转换）。
__double_as_longlong()和__longlong_as-double()强制使寄存器原地配对；__double2float()和__double2hiint()分别返回输入的最低和最高的32位输入操作数；__hiloint2double()则由高位和低位两部分构造一个双精度浮点数。
int double2loint(double d); int double2hint(double d); int hiloint2double(int hi, int lo); double long_as-double(long long int i); long long int __double__as__longlong(double d);8.1.2 本地内存
本地内存是用来容纳寄存器溢出的数据，并存储着被索引的局部变量，这些局部变量的索引在编译时是不能计算的。本地内存同全局内存是由设备内存上同一个内存池所支持的，所以在费米及之后的硬件上，本地内存与L1、L2缓存层次有着同样的延迟特性和优点。本地内存的寻址方式使内存事务可以自动合并（coalesced）。硬件包括了用来加载和存储本地内存的特殊指令：对于特斯拉架构的设备，对应的SASS指令是LLD/LST，而对于费米和开普勒架构的设备是LDL/STL。
8.1.3 全局内存
SM可以用GLD/GST指令（在特斯拉架构设备上）或LD/ST指令（在费米和开普勒架构的设备上）读写全局内存。开发者可以用标准C的操作符来计算和解引用地址，包括指针算术运算和解引用操作符*、[]和->。对64位或128位的内置数据类型（int2,float2/int4,float4）的操作会让编译器自动调用64或128位的加载和存储指令。通过合并内存事务，可以获得最佳的内存性能，详情参见第5.2.9小节。
特斯拉架构的硬件（SM 1.x）使用特殊地址的寄存器来存储指针，之后更新的硬件实现了一个加载和存储架构，对整型与浮点型值
的指针使用同样的寄存器文件，对常量内存、共享内存与全局内存使用相同的地址空间 [1]。
费米架构的硬件包括了许多在更早的设备中不支持的特性
·通过宽加载存储指令支持64位寻址，地址被放在偶数编号的一对寄存器中。在32位的主机平台上不支持64位寻址。在64位主机平台上，64位寻址被自动打开。因此，由同样的内核生成的分别针对32位与64位主机平台的代码可能在寄存器数量与性能上有一些差别。
·一级缓存可以配置成16KB或48KB大小 [2]。（开普勒架构另外可以将缓存分为32KB一级缓存和32KB共享内存。）加载指令可以包含高速缓存提示（来告诉硬件是将数据存入一级缓存，还是越过一级缓存仅将数据保留在二级缓存中）。这些设置可以通过内联的PTX或者命令行选项-X ptxas - dlcm=ca（缓存在一级缓存和二级缓存中，这是默认设置）或者-X ptxas - dlcm=cg（只缓存在二级缓存中）指定。
原子操作（或者仅称为原子）可以在多个GPU线程同时对一个内存位置进行操作时，保证内存更新的正确性。在整个操作期间，硬件会强制在该内存位置上执行互斥访问。因为操作的前后顺序不能保证，所以被支持的操作普遍是符合结合律的 [3]。
在SM 1.1版本中，首次支持了全局内存的原子操作，在更高版本中也同样支持。在SM 1.2及之后的版本中，实现了共享内存的原子操
作。在开普勒架构的设备之前，全局内存的原子操作实际上慢得没有太多使用价值。
在表8-2中总结的全局原子内置函数，当适当的GPU架构通过nvcc选项--gpu-architecture指定后，会自动变得可用。所有的这些函数支持32位整型变量。SM 1.2以后加入了对64位atomicAdd()、atomicExch()和atomicCAS()的支持。SM 2.0中atomicAdd()加入了对32位浮点数（float）的支持，在SM 3.5中加入了atomicMin()、atomicMax()、atomicAnd()、atomicOr()和atomicXor()的64位支持。
表8-2 原子操作
注意 由于原子操作的实现使用了GPU集成内存管理器的硬
件，它们无法穿越PCIe总线工作，所以当设备指针指向主机内存或点对点内存时不能正常工作。
在硬件层面，原子操作分为2种：一种是返回在操作执行前的特定内存位置的值的原子操作；另一种是开发者可以“启动后就不理”的忽略返回值的归约操作。因为在不需要返回值的情况下，硬件可以更
有效地执行操作，编译器会检测返回值是否被使用，如果没有使用的话，发射不同的指令。例如在SM 2.0中，这种指令分别叫ATOM和RED。
8.1.4 常量内存
常量内存驻留在设备内存上，但是，是被另一种只读的缓存支持的，这种缓存经过了优化，可以将读请求的结果广播到所有引用同一内存位置的线程。每个SM包括一个小的、延迟优化的缓存来为那些读请求提供服务。设置内存（和缓存）为只读属性，简化了缓存管理，因为硬件无须实现写回策略来处理被更新的内存。
SM 2.X和后续硬件包括一个对内存进行特别优化的策略，这种优化针对的是并没有被指定为常量内存却被编译器标示为如下两种属性的内存：1)只读；2)地址不依赖线程块或线程的ID。“统一加载”(load uniform, LDU)指令使用常量缓存层次读取内存，并将数据广播到线程中去。
8.1.5 共享内存
共享内存是速度非常快的，是SM的芯片级内存。线程可以使用它在一个线程块中的线程间进行数据交换。每个SM有自己的共享内存，所以共享内存可以影响线程占用率，即一个SM上可常驻的线程束数
量。SM使用特殊的指令来存取共享内存：在SM 1.X上使用G2R/R2G，在SM 2.X及以后的设备上使用LDS/STS。
共享内存由交替排列的存储片（bank）构成，并且通常是针对32位访问来优化的。如果一个线程束中多于一个线程引用同一存储片，一个存储片冲突（bank conflict）就发生了，所以硬件必须连续不断的处理内存请求，直到所有请求被服务。典型的，为了避免存储片冲突，应用程序根据线程编号按照交替模式来访问共享内存，就像下面这样：
extern __shared__float shared[ ]; float data = shared[BaseIndex + threadIdx.x];让线程束中的所有线程从相同的32位共享内存位置读取数据也是非常快的。硬件为处理这种情况，使用一种广播机制。向同一存储片写入的指令将由硬件做序列化处理，降低了性能。向同一个地址的写入操作会引起资源竞争问题，应当被避免。
对2D访问模式（像图像处理内核中的像素分块），最好在分配共享内存时进行补齐，这样内核就可以引用相邻的行，而不至于引起存储片冲突。SM 2.x和之后的硬件拥有32个内存存储片 [4]，所以对2D分块访问，一个线程束中的线程可以按照每一行来访问，补齐分块大小为若干33个32位字是一个很好的策略。
在SM 1.x硬件上，共享内存大约16KB大小 [5]，在之后的硬件上，共有64KB的一级缓存，可以划分出16KB或者48KB的共享内存，剩余的部分继续用作一级缓存 [6]。
在最近的几代硬件中，英伟达公司提升了硬件处理非32位操作数能力。在SM 1.x硬件中，相同存储片的8位和16位操作数的读取会导致存储片冲突，而SM 2.x和之后的硬件可以在相同存储片上广播任意大小的读取。相似的，SM 1.x的共享内存上的64位操作数（像double）操作要远远慢于32位操作数，因此开发者有时不得不把它分割成高32位和低32位来存储。对主要在共享内存使用64位操作数的内核，SM 3.x添加了一个新的特性：一个提升内存存储片大小到64位的模式。
共享内存的原子操作
SM 1.2添加了在共享内存上执行原子操作的功能。不同于全局内存中，使用单指令实现原子操作（根据是否使用返回值分成GATOM或GRED），共享内存原子操作使用显式的加锁/解锁语义实现，编译器生成的代码会在原子操作进行处使每一个线程循环，直到原子操作执行完毕。
代码清单8-1给出了程序atomic32Shared.cu的源代码，意在突出共享内存原子操作中的编译代码生成。代码清单8-2给出了为SM 2.0生成的微码。注意指令LDSLK（加锁加载共享内存）是如何返回一个断定
结果以告诉程序锁是否获得，执行更新的代码进行了分支断定，并且代码将循环等待，直至获得锁并完成更新。
锁在每32位字上执行，锁的索引由共享内存地址的2-9位决定。注意避免资源竞争，否则代码清单8-2中的循环会迭代多达32次。
代码清单8-1 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]; }代码清单8-2 atomic32Shared.cubin（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.CHRED.POPC RZ, RZ;
/*0050*/ LD R0, [R0];
/*0058*/ LDSLK P0, R2, [R3];
/*0060*/ @P0 IADD R2, R2, R0;
/*0068*/ @P0 STSUL [R3], R2;
/*0070*/ @!PO BRA 0x58;
/*0078*/ NOP.S CC.T;
/*0080*/ BAR.CHRED.POPC RZ, RZ;
/*0088*/ LDS R0, [R3];
/*0090*/ IADD R2, R3, c [0x0] [0x24];
/*0098*/ ST [R2], R0;
/*00a0*/ EXIT;8.1.6 栅栏和一致性
我们熟悉的__syncthreads()内置函数会等待线程块中的所有线程到达以后才继续执行。这是维持一个线程块中共享内存一致性所必需的[7]。类似的内存栅栏指令可以被用来在更广范围内的内存上按一定次序执行指令，见表8-3。
表8-3 内存栅栏函数
（续）
[1] 常量与共享内存都存在于地址窗口中，它们即使在64位架构上，也可以被32位地址引用。
[2] 硬件可以在每次内核启动时对此进行配置，但是改变这个状态的代价很大，并且会破坏并发内核启动中的并发性。
[3] 唯一的例外是单精度浮点加法。话又说回来，面对缺少结合律性质的单精度浮点数操作，一般的浮点代码必须是鲁棒的；移植到不同的硬件，甚至只是使用不同的编译器选项重新编译相同的代码，都可以改变单精度浮点数操作的顺序，从而改变结果。
[4] SM 1.x 硬件有着 16 个内存存储片（一个线程束的前 16 个线程到后 16 个线程的内存传输是分别独立服务的），但是在后续硬件上使用的策略对 SM 1.x 同样有效。
[5] 256字节的共享内存被保留，用来进行参数传递；在SM 2.x及其之后，参数改由常量内存传递。
[6] SM 3.x 硬件添加了均分一级缓存的功能, 即 32KB 一级缓存/32KB 共享内存的划分。
[7] 注意：对一个线程束内按“锁步”方式运行的线程，有时允许开发者自己编写“线程束同步”代码，而不用调用__syncthreads()，第7.3小节描述了线程和线程束执行的细节，在本节第三部分中包含了几个不同的线程束同步的代码。