GPU 性能指导： 内存性能

GPU 性能指导： 内存性能
一:内存指令
1：内存指令包括任何从shared，local，globl内存中读或者写指令，仅当存取自动变量时才有可能对local 进行读写；
 2：每个时钟周期可有8个内存操作，但是当存取local和globl 时，还有400个clock cycles的内存延迟作为例子，下面给出一个内存赋值操作的吞吐量：
 _shared__ float shared[32];
_device__ float device[32];
shared[threadidx.x] = device[threadIdx.x]
从global中读一个变量需要8 op，将变量写到shared需要8op，但是从global中读数据还有400-600个时钟周期的延迟；
 3：如果有足够多且独立的运算指令在等待内存读写的时候执行，那么，大多数这种全局内存的延迟可以被线程调度隐藏；

二：同步指令在没有线程等待其它线程的情况下，warp的一条同步指令_syncthreads的性能是每clock 8个操作，也就是4个时钟
三： 内存带宽对各种内存空间带宽的影响直接依赖于内存存取模式，下面几节分别讲解，由于设备内存比芯片内存有更高的延迟和更底的带宽，因此，应该最小化设备内存的存取，典型的编程方法就是将设备内存放到share内存中，也就是说，使一个块的所有线程：
 1：将存储在设备内存中的数据放到共享内存中
 2：为了让每一个线程能安全的读取由不同的线程写入到同一地址的数据，需要同步一个block内其它的线程； 3：在share 内存中处理数据
 4：如果有必要，进行再一次同步，来确保sharememory 中的结果数据被更新
 5：将结果写回到设备内存中四：全局内存全局内存没有缓存，因此，因而采用正确的访问模式来实现最大化的存储器带宽尤为重要，考虑到设备存储器访问的高昂成本时更是如此。首先，设备有能力以一条单一的指令从全局内存中读入一个4,8,16bytes 字到寄存器中：
_device_ type devicep[32];
 type date = device[tid]
将其编译为一条加载指令，type 必须使 sizeof(type) 等于 4、8 或 16，类型为 type 的变量必须与 sizeof(type) 字节一致（也就是说，它的地址必须是 sizeof(type) 的倍数）。对于 4.3.1.1 节介绍的内置类型，如 float2 或 float4，对齐要求可自动满足。对于结构体来说，可通过使用对齐说明符 __align__(8) 或 __align__(16) 来使编译器实施大小和对齐要求，
举例如下：
 struct __align__(8)
{ float a; float b; };
 或：
 struct __align__(16)
{ float a; float b; float c; };
对于超过 16 个字节的结构体来说，编译器会生成多条加载指令。为了确保生成的指令数量最少，应使用 __align__(16) 定义此类结构体，
举例如下：
 struct __align__(16)
 { float a; float b; float c; float d; float e; };
 这将编译为 2 条 128 位的加载指令，而非 5 条 32 位的加载指令。

 2：任何驻留在全局内存或者是被驱动或者运行时函数分配而返回的内存地址，总是至少以256 bytes 对齐；读一个不对齐的8字节或者16字节的字可能导致错误的结果，所有在维护这些类型的任何值或者数组的开始地址时，要特别小心。一种特别的情况就是当数用自定义的数组分配时，使用cudamalloc来一次分配足够的内存在划分给不同的数组使用，可能很容易的就导致内重叠，在这种情况下，每一个数组的开始地址，是block的开始地址的偏移量；
 3：当半warp线程同时存取全局内存时，全局内存的带宽能被更有效的使用，（在执行单一的读或者写指令时）它能将存取合并为单独的32.64或者128 字节的内存传输；
 4：余下的部分描述在各种计算能力的设备上内存联合存储需要满足的条件，如果half-warp满足这些需求，即使warp被分支分开并且一些线程并没有实际的存取内存，内存联合存储也能满足；
 5：为了下面讨论的目的，全局内存被认为是按每段大小为32.64.128划分，并且安这些大小对齐；在计算能力1.0和1.1的设备上合并：如果满足如下三个条件，half-warp的所有线程对全局内存的存取就能合并为1或者2个存取事务.
 1:线程必须存取
 a: 4字节字，结果合并为一个64字节的内存传输
b: 8字节字，结果合并为一个128字节的内存传输
 c: 16字节字，结果合并为两个128字节内存传输
 2：全部的16个字必须位于大小等于内存传输大小的相同段内
3：线程必须依次对这些字进行存取，半Warp块中的第k的线程必须存取第k个字如果半walf块不满足上述需求，那么，将为每个线程发出一个内存传输事务，这将大大的降低性能；实现8字节合并访存比起4字节合并访存来，（a little lower）带宽较底，而实现16字节内存合并访存比起实现4字节来有更底的（lower）。但是，合并访存比起非合并访存来，带宽成数量级的降低，当不合并存取4字节时，比起存取16字节来，底2倍，比起8字节来，底4倍；在计算能力为1.2或着更高的设备上合并访存：一旦half-warp中所有线程存取的字在如下大小的一个段时，那么，所有的half-warp中所有线程对全局内存的访问就能合并到一次单一的内存传输中:
 1：如果所有线程都访问1字节字，则为32字节
 2：如果所有线程都访问2字节字，则为64字节
3：如果所有线程都访问4字节或者8字节，则为128字节； half-warp的任何内存请求模式都将得到合并，包含多个线程存取同一个地址的模式，这和底设备能力中线程存取的字需要是顺序的形成对比；如果 half－warp存取的字在不同的n个内存段中，将有n个内存传输被执行，而对于大于1个内存段的存取，在底计算能力的设备中，将有16次内存传输。特别的，如果线程存取16字节的字，至少要进行两次内存传输；在内存传输中，从不使用的字也被读取，以至于浪费了带宽，为了减少浪费，硬件将自动的调整，仅传输那些被使用的字，例如：如果所有请求的字只有128字节的一半，那么只有64字节被传输：
更为精确的说：下面这些协议用于发布一次half-warp的内存传输：
 1：寻找线程号最小的活动线程所请求的地址所在的内存段，对于1字节数据来说，段大小是32字节，对于2字节数据，则为64字节，对于4或者8字节数据而言，则为128字节。
 2：查找请求地址位于同一段中的其它活动线程
3：降低传输大小
 a:如果传输的内存块大小为128字节并且仅有上一半或者下一半内存被使用，则将传输的大小降到64字节
 b:如果传输的内存块大小为64字节并且仅当上一半（32字节）或者下一半被使用，那么将传输大小减至32字节 4：将传输完成的的线程标记为被服务的不活动线程
 5：重复以上操作，直到half-warp中的所有线程都是以服务线程

 四：共同的存取模式
 1:一个共同的全局内存存取模式是每个线程的线程id存取类型为TYPE ＊基地址位于BaseAddress的数组时，使用下面这种类型地址： BaseAddress + tid
 2:为了实现存储器接合，type 必须满足上文介绍的大小和对齐要求。具体来说，这也就意味着，如果 type 是一个大于 16 字节的结构体，就应将其分割为多个满足这些需求的结构体，而数据应作为此类结构体的多个数组存储于存储器之中，而不是单独一个 type* 类型的数组
3:另外一种全局存储器一般访问模式是索引为 (tx, ty) 的各线程使用以下地址访问类型为 type*、位于地址 BaseAddress 处、宽度为 width 的二维数组中的一个元素： BaseAddress + width * ty + tx 在这种情况下，只有满足了以下条件，才能为线程块的所有半 warp 块实现存储器接合：线程块的宽度是半 warp 块大小的倍数； Width 是 16 的倍数。具体来说，这也就意味着，宽度不是 16 的倍数的数组在分配时的宽度越接近 16 的倍数，并且行得到了相应的填充，那么对于该数组的访问就越有效率。
参考手册中介绍的 cudaMallocPitch() 和 cuMemAllocPitch() 函数和相关的存储器复制函数使程序员能够编写不依赖于硬件的代码，分配符合这些限制条件的数组。
五：局部内存
 1：和全局内存一样，局部内存区也没有cached，所以存取局部内存和存取全局内存的资源消耗一样昂贵，局部内存的存取总是联合的，因为它们是在每个线程中被定义的
2：局部变量存取仅发生在2.5中提及的自动变量时，检查PTX汇编代码(使用-ptx 或者－keep选项）在第一次编译阶段能告诉你一个变量是否真的被放在局部内存区中，它将使用.local助记符标记并使用ld.local和 st.local助记符存取。如果并没有发现，那么可能在随后的编译阶段中，可能将其保存在目标体系的寄存器中，我们不能知道一个变量是被放在寄存器中还是局部内存区中，但是编译器能在使用了--ptxas-option=-v时报告每个内核使用的Lmem总数。
 3：像被放在局部变量中的大结构体或者是数组将耗尽更多的寄存器资源，并且，编译器不能用常数来确定数组的索引；

六：常量内存区常量内存区是有缓存的，所以，仅当cache未命中时才到全局内存中读取，另外，常量内存仅有读操作对于一个half-warp中的所有线程，只要所有的线程都读常量内存的同一地址，那么读取常量缓存和从寄存器中读的速度是一致的。资源消耗随着所有线程读不同的地址而线程增长，我们推荐所有warp的线程都读同一地址，因为可能在以后的设备中，将需要全速读取.

 7：共享存储区
1：因为是芯片级的，shared memory 内存空间的存取比局部和全局内存区快很多。实际上，对于一个warp中的所有线程来说，当存取线程时没有bank conflicts 时，它和读写寄存器一样快.下面详细讨论：

 2：为了实现高内存带宽，share mem 被大小相同的内存模块，称为banks，banks能被同时存取，因此，对在n个独立的bank内存区的n个地址进行读或者写，能被同时服务，这比起在单个bank中读取n次有更高的带宽，并且是单个模块的n倍。然而，如果两个地址的请求同时落在一个bank中，将出现bank conflict 并且存取将被顺序化（一个存取完在进行下一个）硬件将在一个bank 冲突中的请求分离到多个独立的无冲突的请求是必然的，这将有效带宽降低到原来的1/n，如果独立的内存请求是n，在初的内存请求造成n次bank conflicts

 3:为了得到最大的性能，理解为了调度内存请求而如何将内存地址映射到band中是非常重要的，这将导致bank conflicts的最小化。连续的32bit的字被分配给连续的bank，，每个bank的带宽都是每两个时钟周期 32 位。于计算能力为 1.x 的设备，warp 块的大小是 32，存储体的数量为 16（参见第 5.1 节）；
warp 块的共享存储器请求将分割为一个针对 warp 块上半部分的请求和一个针对 warp 块下半部分的请求。因而，属于 warp 块第一部分的线程和属于 warp 块第二部分的线程之间不可能出现存储体冲突。一种常见的情况就是各线程访问数组中的 32 位字，使用线程 ID tid 进行索引，步幅为 s：
__shared__ float shared[32];
float data = shared[BaseIndex + s * tid];
只要s＊n是banks数的倍数或者和banks相等，那么，在这种情况下，线程tid和tid＋n存取同一个bank，每当n是m/d的倍数，而d是m 和s的最大公约数。当warp 的大小的一半小于或者等于m/d时，被认为是没有bank conflict，对于计算能力是 1.x 的设备，可以说只有在 d 等于 1 时，换句话说,只有在 s是奇数时，才不会存在存储体冲突，因为 m 是 2 的幂。其它值得提及的情况是当存取的元素小于或者大于32位时，例如，当以下面这种方式存取一个char型的数组时，就bankconflicts
__shared__ char shared[32];
char data = shared[BaseIndex + tid];
 这是因为shared［0］，shared［1］，shared［2］，shared［3］属性同一bank，如果用下面这种方式存取，则没有bank char data = shared[BaseIndex +4*tid] 下面是double 类型的数组的bank conflicts （第二种方式）因为内存请求被编译成两个独立的32位请求，一种避免这种情况的bank conflicts 是将double操作分离下面这种类型的代码：
 __shared__ int shared_lo[32];
 __shared__ int shared_hi[32];
 double dataIn; shared_lo[BaseIndex + tid] = __double2loint(dataIn);
shared_hi[BaseIndex + tid] = __double2hiint(dataIn); double dataOut =      __hiloint2double(shared_hi[BaseIndex + tid], shared_lo[BaseIndex + tid]);
但这种做法并非总是能够提高性能，在未来的架构中可能表现更差。结构体赋值将在必要时编译为针对结构体中各成员的多个存储器请求，因此，以下代码：
 __shared__ struct type shared[32];
 struct type data = shared[BaseIndex + tid];
 将得到以下结果：
 如果 type 定义如下，则进行三次无存储体冲突的存储器读取：
 struct type { float x, y, z; };
这是因为每个成员都是使用三个 32 位字作为步幅访问的。
  如果 type 定义如下，则进行两次有存储体冲突的存储器读取：
 struct type { float x, y; };
这是因为每个成员都是使用两个 32 位字作为步幅访问的。
  如果 type 定义如下，则进行两次有存储体冲突的存储器读取：
struct type { float f; char c; };
 这是因为每个成员都是使用 5 个字节作为步幅访问的。
最后，当处理一个内存读请求时，shared mem 凭借它的广播机制能在读一个32bit字的同时将它同时广播给几个线程。当一个half-warp的线程从包含同一个32位字的地址读时，这种机制减少了 bank conflicts；更确切的说，一次读请求可能要读多个地址，一个请求由又分几个步骤来完成，每两个clock cycle 一个步骤，每一步通过一个无冲突服务来完成这些地址中的一个子集直到所有的地址都被处理，在每一步中，这个子集都通过下面的几个步骤来构建：
 1：选择指向剩下地址中的一个字作为广播字
 2：在子集中包括
 a：在广播字内的所有地址
 b: 每个bank指向剩余地址的一个指针在每一个周期中，那个字被选择作为广播字以及那个地址放在那个bank中都是未定的，一种常见的conflict -free 情况，就是一个half-warp的所有线程都从包含同一32 bit字的地址处读取数据；