访存优化
CUDA的存储层次包括寄存器、共享内存、本地内存、常量内存、纹理内存、全局内存等，图中描述了CUDA内存空间的层次结构。不同的存储层次存在不同的作用域、生命周期和缓存行为。本节将从全局内存、共享内存、bank冲突以及高速缓存四个方面描述如何对CUDA程序进行访存优化。
全局内存优化
全局内存是GPU中容量最大、延迟最高并且最常使用的存储空间。开发者在主机端可以使用cudaMalloc 函数动态分配全局内存，使用cudaFree函数释放全局内存。在设备代码中使用__device__限定符可以静态的声明一个全局内存变量。在全局内存上分配的空间存在于CUDA程序的整个生命周期中，并且可以被所有核函数中的所有线程访问。因此，若CUDA程序中存在多个核函数使用相同全局变量，则在执行时应密切注意核函数间的内存竞争。
在CUDA执行模型中，对全局内存的访存指令以线程束为单位，并通过缓存来实现加载或存储执行，为了提高CUDA程序对全局内存的访存效率，需要关注两个特性：
（1）合并内存访问，当一个线程束中全部32线程访问一个连续的内存块时，即可达成合并内存访问。
（2）对齐内存访问，当线程束执行内存事务的目标首地址为设备缓存粒度（32字节的二级缓存或128字节的一级缓存）的整数倍时，即可达成对齐内存合并。
在GPU设备上，对全局内存的访问通常需要几百个时钟周期，而执行一次计算操作只需要几个时钟周期，因此除了通过合并对齐提高对全局内存的访问效率之外，提升核函数的计算访存比，复用取自全局内存的数据对提升CUDA程序的性能同样至关重要。
对经过线程结构优化的CUDA矩阵乘的核函数Kernel_2Dgrid2Dblock进行分析，矩阵的数据规模和最优配置线程块的维度均为2的整次幂时，能够实现对全局内存的对齐合并访存。但观察核函数中的核心计算代码sum += Ad[row * width + k] * Bd[k * width + col]可知，块内线程对全局内存进行两次读取操作对应一次乘累加计算操作，计算指令只占计算主体的三分之一，核函数执行中存在大量访问全局内存带来的时延。
为了解决这一问题，重新构建CUDA矩阵乘的核函数，每个线程负责计算一个大小为4×4的矩阵块。经过对CUDA矩阵乘核函数的重新设计，MatrixMul_4x4内计算主体的计算访存比变为了16/8，有利于隐藏访问全局内存时导致的时延。
对经过全局内存优化的CUDA矩阵乘进行测试，在矩阵规模为1024的情况下与经过线程结构优化的CUDA矩阵乘进行对比，编译使用命令为：nvcc global.cu -o global，使用Nsight System工具进行监测核函数运行时间，使用命令：nsys profile –stats=true ./global，测试结果如下表所示。
由测试得到的数据可知，MatrixMul_4x4核函数相较于MatrixMul_2grid2block的执行时间大幅缩小，经过全局内存优化的CUDA矩阵乘实现在((16,16),(16,16))和((8,8),(32,32))的线程布局下性能都远超过10.2节中的CUDA矩阵乘实现，测试结果说明了优化面向全局内存的访问模式能够提升CUDA程序的性能，同时再一次证明了线程布局会影响CUDA程序的性能。
共享内存优化
共享内存是GPU上的关键内存部件，与全局内局相比共享内存具有更高的带宽和更低的延迟，其作用类似于一个可编程管理的缓存，在SM上执行的线程块中的所有线程共享该部分内存空间，因此过度使用共享内存空间会限制SM上活跃线程块的数量。
在CUDA内存模型中，每个线程块在开始执行任务时会被分配一定数量的共享内存空间，该共享内存空间具有与线程块相同的生命周期，且地址空间被线程块中所有的线程共享，因此，共享内存常被用作线程块内线程通信的通道，实现块内线程的相互协同。通过最大化利用共享内存这一高速片上内存资源，可以优化核函数对全局内存访问模式的，提升CUDA程序的性能。
CUDA开发者可以对共享内存变量进行静态或动态的分配，例如一个共享内存的二维浮点数组：__shared__ float tile [size_y][size_x]，使用__shared__修饰符对共享内存变量进行声明，若该共享内存变量在核函数中被声明，则变量的作用域仅为核函数内；若该共享内存变量在CUDA程序中所有核函数外被声明，则变量的作用域应为CUDA程序的全局。
通过进行全局内存优化，CUDA矩阵乘法的性能获得了大幅提升，但全局内存高延迟的物理特性限制了其性能的进一步提升，在此基础上选择grid(8,8)block(32,32)的线程布局，通过使用共享内存资源继续对CUDA矩阵乘进行优化。
首先核函数内使用修饰符__shared__静态开辟了数据规模为1024的共享内存空间ldsa与ldsb，接下来线程根据指令进行数据从全局内存到共享内存的转移，线程块内的1024个线程将矩阵A和矩阵B中的1024（128*8）个元素分别转移至ldsa与ldsb中，线程对结果矩阵块中元素部分和进行计算的核心代码未发生变化，但进行乘累加运算时只需以较低的通信开销到共享内存上获取目标元素，从而大大减少了对全局内存频繁访问带来的时延。
对使用共享内存的CUDA矩阵乘进行测试，编译命令为：nvcc MaMul_shared.cu -o shared，使用Nsight System工具进行监测核函数运行时间，使用命令：nsys profile –stats=true ./shared，测试结果如下表所示。由测试得到的数据可知，核函数MatrixMul_Shared的执行时间仅为255.67us，证明了共享内存的使用成功减少了CUDA矩阵乘中面向全局内存访问带来的时延，CUDA矩阵乘的性能得到了进一步的提升。
由测试得到的数据可知，核函数MatrixMul_Shared的执行时间仅为255.67us，证明了共享内存的使用成功减少了CUDA矩阵乘中面向全局内存访问带来的时延，CUDA矩阵乘的性能得到了进一步的提升。
在以CUDA并行归约例对分支优化进行了说明，CUDA并行归约同样可以利用共享内存提升程序性能，使用共享内存的CUDA并行归约核函数代码如下。
__global__ void reduce_shared(int * g_idata,int *g_odata,){
__shared__ int s_data[1024];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int tx = threadIdx.x;
s_data[cacheIndex] = g_idata[tid];
__syncthreads();
for (int stride = 1; stride < blockDim.x; stride *= 2){
int index = 2 * stride * tx;
if (index < blockDim.x){
s_data[index] += s_data[index + stride];
}
__syncthreads();
}
if (cacheIndex == 0)
g_odata[blockIdx.x] = s_data[tx];
}观察上述代码可知，核函数内开辟了数据规模为1024的共享内存空间s_data，线程块内线程将进行归约计算需要用到的数据块首先从全局内存转移至共享内存上，在进行累加计算时从s_data内获取目标元素，对使用共享内存的CUDA归约进行测试，编译命令为：nvcc Reduce_shared.cu -o Reduce_shared，使用Nsight System工具进行监测核函数运行时间，使用命令：nsys profile –stats=true ./Reduce_shared，测试结果如下表所示。
由测试数据可知，通过与10.3节中CUDA归约的核函数执行时间进行比较，使用共享内存的CUDA归约核函数reduce_shared相较于未经分支优化的CUDA归约核函数reduce_GPU执行时间更短，但性能不及经过分支优化的CUDA归约核函数reduce_NeighboredLess，出现该现象的原因是reduce_shared中对共享内存的访问模式不佳，在下一节中将会针对这一问题展开叙述。
避免bank冲突
内存带宽是衡量存储设备性能的重要指标，为了获得较高的内存带宽，GPU上的共享内存设备被分32个大小相等的存储器模块，这些存储模块被称为存储体bank，可以被一个线程束内的32个线程同时访问。如下图所示，在费米架构的GPU设备上，连续的4字节数据被分配到连续的32个存储体中。在费米架构的GPU设备中存储体的宽度为4字节，在开普勒及之后架构的GPU设备中存储体的宽度为8字节。
当一个线程束中的不同线程访问一个bank中的不同的字地址时，就会发生bank冲突。图中展示了三种不同的共享内存访问模式。
线性访问模式中，线程束内的线程步长为1，访问过程中线程ID与存储体ID一一对应。交叉访问模式中线程ID虽然没有与存储体ID一一对应，但线程束内的每个线程对应一个唯一的存储体。以上两种访问模式中均不存在bank冲突，单次访问操作可以由一个内存事务实现。不规则访问模式中，若线程束内的多个线程访问同一个存储体中的相同地址，则该模式被称为多播模式，通过广播访问可以避免bank冲突的发生。若线程束内的多个线程访问同一个存储体内的不同地址，则会出现bank冲突。在使用共享内存时若访问模式中出现了bank冲突，会降低对共享内存的访问效率，在最不理想模式下，若一个线程束中的所有线程访问相同存储体中的32个不同字地址，则该访问操作需要由32个内存事务完成，即产生了32路bank冲突，从而严重降低了内存带宽。
对bank冲突的相关概念进行了解后，重新分析10.4.3节中使用共享内存的CUDA归约核函数代码，发现执行累加操作的s_data[index] += s_data[index + stride]语句会在读取s_data内目标元素时导致bank冲突，当步长变量stride为1时，一个线程束对s_data的访问会产生两路bank冲突，随着迭代中步长的变量stride的增长，bank冲突现象更加严重。通过重新构建累加操作的执行方式来避免bank冲突。
对经过bank冲突优化的共享内存CUDA归约进行测试，编译使用命令：nvcc bankconflict.cu -o bankconflict，使用Nsight System工具进行监测核函数运行时间，使用命令：nsys profile –stats=true ./ bankconflict，测试结果如下表所示。
使用Nsight Compute工具对共享内存事务进行监测，其中l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum与l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum选项分别表示核函数执行过程中对共享内存进行读写所需内存事务的总和，smsp__sass_average_data_bytes_per_wavefront_mem_shared.pct参数表示核函数对共享内存的利用效率，使用命令ncu –metrics smsp__sass_average_data_bytes_per_wavefront_mem_shared.pct ./bankconflict，测试结果如下表所示。
高速缓存优化
与CPU缓存类似，GPU缓存是不可被编程的内存空间。在GPU上有4种缓存分别为一级缓存、二级缓存、只读常量缓存以及只读纹理缓存。在每SM中有一个只读常量缓存和只读纹理缓存，它们用于进一步提高GPU设备的读取性能。在CPU上，内存数据的加载和存储都可以被缓存，但是GPU上的内存存储操作不会被缓存，只有内存加载操作会被缓存。
下面以CUDA矩阵转置为例说明利用高速缓存的优化，将矩阵A、矩阵B均存在全局内存中，其中核函数transpose1中按行对矩阵A进行合并的读操作，而对矩阵B的写操作是非合并的。在核函数transpose2中按列对矩阵A进行非合并的读操作，对矩阵B进行合并的写操作。
对两种CUDA矩阵转置实现进行测试，编译使用命令：nvcc cache.cu -o cache，使用Nsight System工具进行监测核函数运行时间，使用命令：nsys profile –stats=true ./cache。测试结果如下表所示。
由测试结果得到的数据可知，transpose2的执行时间远小于transpose1，出现性能差距的原因是transpose2对矩阵A的非合并读操作会经由高速缓存，而transpose1中对矩阵B的非合并写操作并不能被缓存，transpose2利用高速缓存优化了面向全局内存的不合并访问，从而获得了更优的性能。
一级缓存和共享内存共享SM上的内存资源，可以通过cudaFuncSetCacheConfig API动态的分配二者的资源占比，其函数原型如下：
cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCachecacheConfig);
func表示分配策略：
cudaFuncCachePreferNone: no preference (default)
cudaFuncCachePreferShared: prefer 48KB shared memory and 16KB L1 cache
cudaFuncCachePreferL1: prefer 48KB L1 cache and 16KB shared memory
cudaFuncCachePreferEqual: Prefer equal size of L1 cache and shared memory, both 32KB优化人员可以通过调用cudaFuncSetCacheConfig函数并选用适当的分配策略对GPU设备上一级缓存与共享内存资源的比例进行调整，从而实现对CUDA程序的优化。
