CUDA学习之第四章：全局内存（一）

4.1 CUDA内存模型概述

4.1.1 内存层次结构的优点

应用程序有两种局部性：时间局部性和空间局部性。

时间局部性：一个数据被引用后，较短时间内会再用，时间越久越不太可能用。
空间局部性：一个数据被引用后，附近的数据会被引用。

现代的计算机支持对局部性好的程序优化性能就更好。对于内存层次结构来说，越上层速度越快，容量越小，越下层速度越慢，但容量越大。
对于CUDA来说，可以显式控制内存层次的行为。

4.1.2 CUDA内存模型

CUDA有六种可编程内存类型：寄存器、共享内存、本地内存、常量内存、纹理内存、全局内存。

其中全局内存、常量内存、纹理内存是与主机交互的，寄存器、共享内存、本地内存是块内的。
一个核函数中的线程都有自己的本地内存，一个块有自己的共享内存，对于其中之线程都是可见的。其内容与线程块一起存在，所有线程都可以访问全局内存、常量内存、纹理内存，但是后两者是只读的。纹理内存是为各种数据布局提供了不同的寻址模式和滤波模式。

4.1.2.1 寄存器

寄存器是最快的内存空间，在核函数中没有声明其他修饰符的自变量通常存储在寄存器中。对于声明的数组，如果其索引编译时可以确定那么也在寄存器中。
寄存器是由活跃线程束划分出的资源，少用寄存器可以使SM上有更多的常驻线程块。
用nvcc编译器选项可以检查核函数硬件资源情况，下面命令可以查寄存器数量、共享内存字节数、每个线程所使用的常量内存的字节数。

-Xptxas -v,-abi=no

（不知道咋用这个命令，我在编译文件后添加这条会报错，如果不加逗号和后面的-abi=no，可以得到寄存器数量和常量内存字节数）
如果超出硬件限制数量的寄存器，会用本地内存替代多占用的寄存器，影响性能。
在代码中可以显式地加上额外的信息帮助编译器进行优化：

__global__ void
__launch_bounds__(maxThreadPerBlock,minBlocksPerMultiprocessor)
kernel(...)
{
	...
}

其中maxThreadPerBlock指出了每个线程块可包含的最大线程数，minBlocksPerMultiprocessor是可选参数，指定每个SM预期最小常驻线程块数量。
还可以使用-maxrregcount编译器选项，控制编译单元里所有核函数使用寄存器的最大数量：

-maxrregcount=32

使用了指定的启动边界时此处指定的值（32）会失效。

4.1.2.2 本地内存

上文已提到寄存器溢出之变量会存于本地内存。编译器可能存放到本地内存中的变量有：

• 编译时使用未知索引引用的本地数组
• 可能占用大量寄存器的本地结构体和数组
• 不满足寄存器限定条件的变量

本地内存与全局内存放在一块的，对于计算能力2.0的GPU，本地内存数据放在SM的一级缓存和每个设备的二级缓存中。

4.1.2.3 共享内存

使用以下修饰符之变量将存在共享内存中：

__shared__

共享内存是片上的，所以飞快，使用是类似于CPU的一级缓存的，但是可编程。
生命周期和块相同，块结束后会释放。
使用时也不能用太多，否则会限制活跃线程数数量。一个块中线程可以通过共享内存通信，访问是需要用同步：

void __syncthreads();

此函数设立障碍点，使得所有线程必须在其他线程被允许执行前到达该处。避免数据冲突。当然这个函数会影响效率，他会强制SM处于空闲状态。
SM中的一级缓存和共享内存都使用64KB的片上内存。它们通过静态划分。但是运行时可以通过如下指令进行动态配置：

cudaError_t cudaFuncSetCacheConfig(const void* func,enum cudaFuncCacheConfig cacheconfig)

此函数会用cudaFuncCacheConfig对于核函数func指定片上内存划分。四种配置如下：

 - cudaFuncCachePreferNone：没有参考值（默认）
 - cudaFuncCachePreferShared：建议48KB共享和16KB一级缓存
 - cudaFuncCachePreferL1：建议48KB一级缓存和16KB共享
 - cudaFuncCachePreferEqual：建议都是32KB

Fermi设备不支持第四种配置。

4.1.2.4 常量内存

常量内存驻留与设备内存，在每个SM中有专用的常量缓存中给他缓存。用以下修饰符修饰：

__constant__

常量变量必须在全局空间内和所有核函数外声明，且所有计算能力都只能声明64KB的常量内存。它是静态声明的，并对同一编译单元所有核函数可见。
核函数只能读取之，故必须于主机端用此函数来初始化常量内存。

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);

它将count个字节从src复制到symbol，此函数大部分情况是同步的。
当所有线程都要读取同一个内存地址之数据时，常量内存效果最好，因为他是广播形式的，如果每个线程都不一样就别用他。

4.1.2.5 纹理内存

在SM中的只读缓存中缓存，纹理内存是通过指定只读缓存来访问的全局内存，其包括对硬件滤波的支持，访问二维数据可以达到最优性能。对于某些程序有奇效，但是另一些程序还不如全局内存。

4.1.2.6 全局内存

这是GPU中最大，延迟最高的内存，可以在任何SM设备上被访问，贯穿应用程序整个生命周期。可以静态或动态声明，用以下修饰符在设备代码中声明变量：

__device__

在第二章中的cudaMalloc和cudaFree就是动态分配全局内存。还是老问题，多个线程访问时要注意，由于不能同步，并发修改同一位置的全局内存可能会导致出问题。
全局内存常驻于设备内存，可以通过32字节、64字节或128字节的内存事务进行访问，这些内存事务必须自然对其，就是说它首地址必须是32、64或128字节的倍数。优化内存事务对于性能很关键。
（一次内存事务表示将32字节数据从全局内存传输到SM，对于一个线程束访问4字节的单精度浮点，那么就是请求128字节的数据）
当一个线程束执行内存加载或存储时，需要满足的传输数量通产取决于以下因素：

• 跨线程的内存地址分布
• 每个事务内存地址的对齐方式

一般来说用来满足内存请求的事务越多，未使用的字节被传回去的可能性越大，这会影响吞吐率。缓存的内存事务利用数据局部性来提升吞吐率。

4.1.2.7 GPU缓存

GPU缓存是不可编程的，有四种：

• 一级缓存
• 二级缓存
• 只读常量缓存
• 只读纹理缓存

每个SM都有一个一级缓存，所有SM共享一个二级缓存，都用来存储本地内存或全局内存的数据，以及寄存器溢出的部分。
GPU中对于内存加载操作可以缓存，而内存存储操作不能。
每个SM中只有一个只读常量缓存和只读纹理缓存。

4.1.2.8 CUDA变量声明总结

就是个图，总结了一下，跳过了。

4.1.2.9 静态全局内存

书中给了个静态使用全局内存的例子，一个核函数中用了个全局内存一直累加浮点数，注意几点：

• 定义在了全局变量的位置（也就是所有函数外面，一般写在预编译下面），并使用修饰符，具体为：__device__ float devData
• 使用cudaMemcpyToSymbol(devData,&value,sizeof(flaot))来传递初始值，value是一个float浮点数。核函数运行完传回的时候使用cudaMemcpyFromSymbol(&value,devData,sizeof(float))。
• 在上述函数中，主机端其实无法访问设备，是API偷偷用了GPU硬件访问的。此处的devData就是个标识，他不是全局内存的变量地址。在核函数中devData是全局内存的一个变量。
• 此处绝不能用cudaMemcpy来传递，因为devData根本不是全局内存的真地址。如果非要用，需要先用这个函数：cudaError_t cudaGetSymbolAddress(void** devPtr,const void* symbol);来获取devData真实地址，然后再用cudaMemcpy。（cuda固定内存是可以直接在主机端引用的，之后会学习这个。）