??大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。因此，最大限度地利用全局内存带宽是调控核函数性能的基础。如果不能正确的调控全局内存的使用，其他优化方案很可能也收效甚微。为了在读写数据时达到最佳的性能，内存访问操作必须满足一定的条件。CUDA执行模型的显著特征之一就是指令必须以线程束为单位进行发布和执行。存储操作也是同样。在执行内存指令时，每个线程都提出了一个包含请求地址的单一内存访问请求，它并由一个或多个设备内存传输提供服务。根据线程束中内存地址的分布，内存访问可以被分成不同的模式。接下来开始介绍如何实现最佳的全局内存访问。
对 齐 与 合 并 访 问
??如下图所示，全局内存通过缓存来实现加载/存储。全局内存是一个逻辑内存空间，我们可以通过核函数访问它。所有的应用程序数据最初存在于DRAM上，即物理设备内存中。核函数的内存请求通常是在DRAM设备和片上内存间以128字节或32字节内存事务来实现的。所有对全局内存的访问都会通过二级缓存，也有很多访问会通过一级缓存，这取决于访问类型和GPU架构。如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现的。对全局内存缓存其架构，如果允许使用一级缓存，那么可以在编译时选择启用或禁用一级缓存。一行一级缓存是128个字节，他映射到设备内存上一个128字节的对齐段。如果线程束的每个线程请求一个4字节的值，那么每次请求就会获得128字节的数据，这恰好与缓存行和设备内存段的大小相契合。因此在优化应用程序时，我们需要注意设备内存访问的两个特性：对齐内存访问；合并对齐访问。
??当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时(32字节的二级缓存或128字节的一级缓存)，就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。对齐合并内存访问的理想状态是线程束从对齐内存地址开始访问一个连续的内存块。为了最大化全局内存吞吐量，为了组织内存操作进行对齐合并是很重要的。下图描述了对齐与合并内存的加载操作。在这种情况下，只需要一个128字节的内存事务从设备内存中读取数据。
??下图扎实了非对齐和未合并的内存访问，在这种情况下，可能需要3个128字节的内存事务来从设备内存中读取数据：一个在偏移量为0的地方开始，读取连续地址之后的数据，一个在偏移量为256的地方开始，读取连续地址之前的数据；另一个在偏移量为128的地方开始读取大量的数据。注意在内存事务之前和之后获取的大部分字节将不能被使用，这样会造成带宽浪费。
??一般来说，需要优化内存事务效率：用最少的事务次数满足最多的内存请求。事务数量和吞吐量的需求随设备的计算能力变化。
全 局 内 存 读 取
??在SM中，数据通过以下3种缓存/缓冲路径进行传输，具体使用何种方式取决于引用了哪些类型的设备内存：一级和二级缓存、常量缓存、只读缓存。其中一/二级缓存是默认路径。想要通过其他两种路径传输数据需要应用程序显示地说明，但要想提升性能还是取决于使用的访问模式。全局内存加载操作是否会通过一级缓存取决于两个因素：设备的计算能力、编译器选项。内存加载可以分为两类：缓存加载（启用一级缓存）、没有缓存的加载（禁用一级缓存）。内存加载的访问模式有如下特点：1.有缓存与没有缓存：如果启用一级缓存，则内存加载被缓存；2.对齐与非对齐：如果内存访问的第一个地址是32字节的倍数，则对齐加载；3.合并与非合并：如果线程束访问一个连续的数据块，则加载合并。我们可以通过编译器标志启用或禁用全局内存负载的一级缓存，如果一级缓存被禁用，所有对全局呢村的加载请求将直接进入到二级缓存；如果二级缓存缺失，则由DRAM完成请求。
缓存加载
??缓存加载经过一级缓存。在粒度为128字节的一级缓存上由设备内存事务进行传输。缓存加载可以分为对齐/非对齐及合并/非合并。下图是理想情况的对齐与合并内存访问。线程束种所有线程请求的地址都在128字节的缓存行范围内。完成内存加载操作只需要一个128字节的事务。总线的使用率为100%，在这个事务中没有未使用的数据。
??下图所示另一种情况：访问是对其的，引用的地址不是连续的线程ID，而是128字节范围内的随机值。由于线程束中线程请求的地址仍然在一个缓存行范围内，所以只需要一个128字节的事务来完成这一内存加载操作。总线利用率仍然是100%，并且只有当每个线程请求在128字节范围内有4个不同的字节时，这个事务中才没有未使用的数据。
??下图也说明了一种情况：线程束请求32个连续4个字节的非对齐数据元素。在全局内存中线程束的线程请求的地址落在两个128字节段范围内。因为当启用一级缓存时，由SM执行的物理加载操作必须在128个字节的界线上对齐，所以要求有两个128字节的事务来执行这段内存加载操作。总线利用率50%,并且在这两个事务中加载的字节有一半是未使用的。
??下图说明了一种情况：线程束中所有线程请求相同的地址。因为被引用的字节落在一个缓存行范围内，所以只需要请求一个内存事务，但总线利用率非常低。如果加载的值是4字节的，则总线的利用率是4字节请求/128字节加载 = 3.125%。
??下图为最差的情况：线程束中线程请求分散于全局内存中的32个4字节地址。尽管线程束请求的字节总数仅为128个字节，但地址要占用N个缓存行(0<N<=32)。完成一次内存加载操作需要申请N次内存事务。
??CPU一级缓存优化了时间和空间局部性。GOU一级缓存是专为空间局部性不足而不是为了时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。
没有缓存的加载
??没有缓存的加载不经过一级缓存，它在内存段的粒度上(32个字节)而非缓存池的粒度(128个字节)执行。这是更细粒度的加载，可以为非对齐或非合并的内存访问带来更好的总线利用率。
??下图所示为理想情况:对齐与合并内存访问。128个字节请求的地址占用了4个内存段，总线利用率为100%。
??下图说明了一种情况：内存访问是对齐的且线程访问是不连续的，而是在128个字节范围内随机进行。只要每个线程请求唯一的地址，那么地址将占用4个内存段，并且不会有加载浪费。这样的随机访问不会抑制内核性能。
??下图说明了一种情况：线程束请求32个连续的4字节元素但加载没有对齐到128个字节的边界。请求的地址最多落在5个内存段内，总线利用率至少为80%，与这些类型的请求缓存加载相比，使用非缓存加载会提升性能，这是因为加载了更少的未请求字节。
??下图说明了一种情况：线程束种所有线程请求相同的数据。地址落在一个内存段内，总线的利用率是请求的4字节/加载的32字节=12.5%，在这种情况下，非缓存加载性能也是优于缓存加载的性能。
??下图说明了最坏的一种情况:线程束请求32个分散在全局内存中的4字节地址。由于请求的128个字节最多落在N个32字节的内存分段内而不是N个128字节的缓存行内，所以相比于缓存加载，即使是最坏的情况也有所改善。
非对齐读取的示例
??因为内存模式往往是由应用程序实现的一个算法来决定的，所以对于某些应用程序来说合并内存加载是一个挑战。然而，在大多数情况下，使用某些方法可以帮助对齐应用程序内存访问。为了说明核函数中非对齐访问对性能的影响，我们对向量代码进行修改，去掉所有的内存加载操作，来指定一个偏移量。注意在下面的核函数中使用了两种索引，新的索引k由给定的偏移量上移，由于偏移量的值可能会导致加载出现非对齐加载。只有加载数组A和数据B的操作会用到索引k。对数组C的写操作仍操作仍使用原来的索引i，以确保写入访问保持对齐。
__global__ void readOffset(float *A, float *B, float *C, const int n, int offset)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int k = i + offset;
if(k < n)
{
C[i] = A[i] + B[i];
}
}
??为保证修改后核函数的正确性，主机代码也应该做出相应的修改：
void sumArrayOnHost(float *A, float *B, float *C, const int n, int offset)
{
for(int idx = offset, k = 0; idx < n; idx++, k++)
{
C[k] = A[k] + B[k];
}
}
只读缓存
??只读内存最初是预留给纹理内存加载使用的，对计算能力为3.5及以上的GPU来说，只读缓存也支持使用全局内存加载代替一级缓存。只读缓存的加载粒度是32个字节。通常，对分散读取来说，这些更细粒度的加载要优于一级缓存。有两种方式是可以直到内存通过只读缓存进行读取：1.使用函数__Idg；2.在间接引用的指针上使用修饰符。例如，考虑下面的拷贝核函数,可以用内部函数__Idg来通过只读缓存直接对数据进行读取访问。
__global__ void copyKernel(int *out,int *in)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
out[idx] = in[idx];
}
__global__ void copyKernel(int *out,int *in)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
out[idx] = __Idg(&in[idx]);
}
??我们可以将__restrict__修饰符应用到指针上，这些修饰符帮助nvcc编译器识别无别名指针(即专门用来访问特定数组的指针)。nvcc将自动通过只读缓存指导无别名指针的加载。
__global__ void copyKernel(int * __restrict__ out,int * __restrict__ in)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
out[idx] = in[idx];
}
全 局 内 存 写 入
??内存的存储操作相对简单，一级缓存不能用在Fermi或Kepler GPU上进行存储操作，在发射到设备内存之前存储操作只通过二级缓存。存储操作在32个字节段的粒度上被执行。内存事务可以被同时分为一段、两段或四段。例如，如果两个地址同属于一个128个字节区域，但是不属于一个对其的64字节区域，则会执行一个四段事务（也就是说，执行一个四段事务比执行两个一段事务效果更好）。
??下图所示为理想情况：内存访问是对齐的，并且线程束里所有的线程访问一个连续的128字节。存储请求由一个四段事务完成。
??下图所示为内存访问是对齐的，但地址分散在一个192字节范围内的情况，存储请求由3个一段事务来实现。
??下图所示的内存访问是对齐的，并且地址访问在一个来纳许的64个字节范围内的情况，这种存储请求由一个两段事务来完成。
??为了验证非对齐对内存存储效率的影响，按照下面的方式修改向量加法核函数，仍然使用两个不同的索引：索引k根据给定的偏移量进行变化，而索引i不变(并因此产生对齐访问)。使用对齐索引i从数组A和数组B中进行加载，以产生良好的内存加载效率。使用偏移量索引x写入数组C，可能会造成非对齐写入，这取决于偏移量的值。
__global__ void writeOffset(float *A, float *B, float *C, const int n, int offset)
{
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int k = i + offset;
if(k < n)
{
C[i] = A[i] + B[i];
}
}
void sumArrayOnHost(float *A, float *B, float *C, const int n, int offset)
{
for(int idx = offset, k = 0; idx < n; idx++, k++)
{
C[k] = A[k] + B[k];
}
}
结 构 体 数 组 与 数 组 结 构 体
??下面通过代码来区分一下数组结构体(AoS)和结构体数组(SoA)：
struct innerStruct{
float x;
float y;
}
struct innerStruct myAos[N]; //数组结构体
struct innerArray{
float x[N];
float y[N];
}
struct innerArray moa; //结构体数组
??下图说明了AoS和SoA方法的内存布局。用AoS模式在GPU上存储示例数据并执行一个只有x字段的应用程序，将导致50%的带宽损失，因为y值在每32个字节段或128个字节缓存行上被隐式地加载。AoS格式也在不需要的y值上浪费了二级缓存空间。用SoA模式存储数据充分利用了GPU地内存带宽。由于没有相同字段元素的交叉存取，GPU上的SoA布局提供了合并内存访问，并且可以对全局内存实现更高效的利用。
??许多并行编程范式，尤其是SIMD型范式，更倾向于使用SoA。在CUDA C编程中也普遍倾向于使用SoA,因为数据元素是为全局内存的有效合并访问而预先准备好的，而被相同内存操作引用的同字段数据元素在存储时是彼此相邻的。
性 能 调 整
??优化设备内存带宽利用率有两个目标：1.对齐及合并内存访问，以减少带宽的浪费；2.足够的并发内存操作，以隐藏内存延迟。影响设备内存操作性能的因素主要有两个：1.有效利用设备DRAM和SM片上内存之间的字节移动：为了避免设备内存带宽的浪费，内存访问模式应该是对齐和合并的；2.当前的并发内存操作数：可通过以下两点实现最大化当前存储器操作数。1）展开，每个线程产生更多的独立内存访问，2）修该核函数启动的执行配置来使每个SM有更多的并行性。