在第六章中，抽象地介绍了CUDA中的各种内存，从本章开始，将通过实例讲解各种内存的合理使用。在各种设备内存中，全局内存具有最低的访问速度(最高的延迟)，往往作为一个CUDA程序的门槛，所以值得特别的关注。本章讨论全局内存的合理使用。

7.1 全局内存的合并与非合并访问

        对全局内存的访问将触发内存事务(memory transaction),也就是数据传输(data transfer).从费米架构开始有了SM层面的L1缓存和设备层面的L2缓存，可以用于缓存全局内存的访问。在启用了L1缓存的情况下，对全局内存的读取首先尝试通过L1缓存；如果未命中，则接着尝试读取L2缓存；如果再次未命中，则直接从DRAM读取。一次数据传输的数据量默认情况下是32字节。
        关于全局内存的访问模式，由合并(coalesced)与非合并(uncoalesed)之分，合并访问是指一个线程束对全局内存的一次访问请求(读或者写)导致最少数量的数据传输，否则称访问是非合并的。定量的说，可以定义一个合并度(degree of coalescing),它等于线程束请求的字节数除以由该请求导致的所有数据传输处理的字节数。若所有数据传输中处理的数据都是线程束所需要的，那么合并度就是100%，即对应合并访问。所有可以将合并度理解为一种资源利用率。利用率高，核函数中与全局内存访问相关的部分性能就更好；利用率低则意味着对于显存带宽的浪费。
        本章主要以全局内存的读取和仅利用L2缓存的情况为例进行讨论。在此情况下，一次数据传输指的就是将32字节的数据从全局内存(DRAM)通过32字节的L2缓存片段传输到SM。考虑一个线程束访问单精度浮点型的全局内存变量的情形，因为一个单精度浮点数占有4个字节，所以该线程束将请求128字节的数据。在理想情况下(即合并度100%)这将触发128/3=4次用L2缓存的数据传输。那么在什么条件下会导致多于四次的数据传输呢。
        为了回答这个问题，首先需要了解数据传输对数据地址的要求：在一次数据传输中，从全局内存转移到L2缓存的一片内存的首地址一定是最小粒度(32字节)的整数倍。例如，一次数据传输只能从全局内存读取地址为0-31字节、32-63字节、64-95字节、96-127字节等片段的数据。如果线程束，如果线程束请求的全局内存数据的地址刚好是0-127或者128-255字节等，就能与四次传输所处理的数据完全吻合。这种情况下的访问就是合并访问。
        读者也许会问：如何保证一次数据传输中内存片段的首地址为最小粒度的整数倍，或者：如何控制所使用的全局内存的地址。答案是：使用CUDA运行时函数的API函数(cumaMalloc())分配的内存的首地址至少是256字节的整数倍。
        下面将通过几个具体的核函数列举几种常见的内存访问模式和合并度。
        (1)顺序的合并访问。我们考察如下的核函数和相应的调用：
 

        void __global__ add(float *x,float *y,float *z)
{
        int n=threadIdx.x+blockDim.x*blockIdx.x;
        z[n]=x[n]+y[n];
}
add<<<128,32>>>(x,y,z);

其中，x,y,z是有cudaMalloc()分配的全局内存的指针。例如第一个线程块中线程束将访问数组x中第0-31个元素，对应128字节的连续内存，而且首地址一定是256字节的整数倍，这样访问只需要4次数据传输就可以完成，所以是合并访问，合并度为100%。
        (2)乱序的合并访问。将上述函数稍稍修改：

void __global__ add_permuted(float *x,float *y,float *z)
{
    int tid_permuted=threadIdx.x^0x1;
    int n=tid_permuted+blockIdx.x*blockDim.x;
    z[n]=x[n]+y[n];

    
}
add_permuted<<<128,32>>>(x,y,z);

其中，threadIdx.x^0x1是某种置换操作，作用是将0-31的整数做某种置换(交换两个相邻的数)。第一个线程块中的线程束将依然访问是访问数组x的0-31号元素，只不过线程号与数组元素指标不完全一致而已，这种访问是乱序的，合并度也是100%。
        (3)不对齐的非合并访问，将第一个核函数稍作修改：

void __global__ add_offset(float *x,float *y,float*z)
{
    int n=threadIdx.x+blockDim.x*blockIdx.x+1;
    z[n]=x[n]+y[n];
    
}
add_offset<<<128,32>>>(x,y,z)

第一个线程块中将访问数组x中第1-32个元素，加入元素的首地址为256字节，该线程束将访问设备内存的260-387字节，这将触发五次数据传输，对应的内存地址为256-287,288-319,320-351,352-383和384-415，这样的访问属于不对齐，合并度为4/5=80%。
        (4)跨越式的非合并访问
将第一个核函数改写如下：

void __global__ add_stride(float *x,float *y,float *z)
{
    int n=blockIdx.x+threadIdx.x*grimDim.x;
    z[n]=x[n]+y[n];
}
add_tride<<<128,32>>>(x,y,z);

第一个线程块中的线程束将访问数组x中指标为0,128,256,384等元素，因为这里的每一对元素都不在一个连续的32字节的内存片段，所以将触发32次数据传输，这样的访问属于跨域式的非合并访问，合并度为4/32=12.5%.
        (5)广播式的非合并访问
 

void __global__ add_broadcast(float *x,float *y,float *z)
{
    int x=threadIdx.x+blockIdx.x*blockDim.x;
    z[n]=x[0]+y[n];
}
add_broadcast<<<128,32>>>(x,y,z);

第一个线程块中的线程束将一致地访问数组x中的第0个元素。这只需要一次数据传输(处理32个字节)，但由于整个线程束只使用了四字节的数据，所以合并度为4/32=12.5%。这样的访问属于广播式的非合并访问，这样的访问适合采用常量内存。

7.2 例子：矩阵转置

        本节将通过一个矩阵转置的力磁讨论全局内存的合理使用。矩阵转置是线性代数中的常规操作。这里仅考虑行数和列数相等的方阵。
        假设一个矩阵A的矩阵元为Aij，则其转置矩阵B=AT的，例如
 则其转置为

7.2.1 矩阵复制

        在讨论矩阵转置之前，先考虑一个更简单的问题：矩阵复制，即形如B=A的计算，以下代码给出了矩阵复制核函数copy()的定义
 

//设计处理的过程中是128*128的矩阵，线程块大小为32*32，从而N取为128，
//即将要处理矩阵的维度
const int Nx=128;
const int Ny=128;
const int block_size_x=32;
const int block_size_y=32;
const int grid_size_x=(N+block_size_x-1)/block_size_x;
const int grid_size_y=(N+bclok_size_y-1)/bclock_size_y;
const dim3 block_size(block_size_x,block_size_y);
const dim3 grid_size(grid_size_x,grid_size_y);
copy<<<grid_size,block_size>>>(d_A,d_B)

__global void copy(const real *A,const real *B)
{
    const int nx=threadIdx.x+blockDim.x*blockIdx.x;
    const int ny=threadIdx.y+blockDim.y*blockIdx.y;
    const int index=ny*blockDim.x*gridDim.x+nx;
    if(nx<blockDim.x*gridDim.x&&ny<blockDim.y*gridDim.y)
        {
            B[index]=a[index];

        }
    
}

首先说明一下，在核函数中可以直接使用函数外部由#define或const定义的常量,包括浮点数和整形常量，但微软的编译器MSVC使用时有一个限制，即不能在核函数中使用在函数外部由const定义的浮点型常量，但在本例中block_dim_x和block_dim_y都是整形常量在文件开头进行定义，也可以改变成#define，可以在核函数中直接进行操作，但是不能使用常量的应用和地址。
        在看copy()的执行配置。在调用copy()的时候，用了二维网格和二维线程块，在该问题中并不是一定要应用二维网格和线程块，因为矩阵的数据排列本质上是一样的。然而在矩阵的转置问题中，使用二维网格和线程块更为方便。如上所述，程序中block_size_x和block_size_y为整形常量，指的是一片(tile)矩阵的维度(dimension,即行数和列数)。将通过线程块一片一片的处理一个大矩阵(128*128)。其中一片矩阵是32*32.每一个二维的线程块处理一片矩阵。线程块的维度和一片矩阵的维度一致。和线程块一致，网格也用二维的，维度为待处理矩阵的维度除以线程块维度。例如假设N=128，则grid_size_x和grid_size_y都是4.也就是说核函数所用的网格维度为4*4，线程块维度为32*32。此时在核函数的blockDim.x和blockDim.y都等于32，且一个线程块中的线程数目为1024刚好等于所允许的最大值。
        最后看核函数最后的实现，nx为矩阵的列标与带.x的内建变量联系起来，ny为矩阵的行标，将其与.y的内建变量联系起来。在其之后将上述行指标和列指标结合起来转化为一维指标index，并且在不越界的条件下将A的第indx元素赋值给B的index元素。
        继续来分析一下核函数中对全局内存的访问模式，对于多维线程块而言，x维度的线程指标threadIdx.x在最内层(变化最快)，所以相邻的threadIdx.x对应相邻的线程。从核函数中的代码可知，相邻的nx对应相邻的线程，也对应相邻的数组元素(对A和B都成立)。所以，在核函数中，相邻的线程访问了相邻的数组元素，没有内存不对齐的情况出现。

7.2.2 使用全局内存进行矩阵转置

        在7.2.1节中讨论了矩阵复制的计算，本节将讨论矩阵转置的计算，为此回顾一下7.2.1节矩阵复制核函数中的如下语句：
                const int index=nx+ny*gridDim.x*blockDim.x;
                if(nx<gridDim.x*blockDim.x&&ny<gridDim.y*blockDim.y) B[index]=A[index]
为了便于理解，我们首先将这两条语句写成一条语句；
if(nx<gridDim.x*blockDim.x&&ny<gridDim.y*blockDim.y) B[nx+ny*gridDim.x*blockDim.x]=A[nx+ny*gridDim.x*blockDim.x]：
在数学的角度上相当于做了Bij=Aij的操作，若要实现矩阵转置即Bij=Aij可以将上述代码改变为
B[nx*gridDim.y*blockDim.y+ny]=A[ny*gridDim.x*blockDim.x+ny];
B[ny*gridDim.x*blockDim.x+nx]=A[nx*gridDim.y*blockDim.y+ny]；
以上语句均能实现矩阵转置，可以看出在第一种方式中对于矩阵A的访问(读取)是顺序的，但是对矩阵B中的访问(写入)不是顺序的。但对于第二种方式而言A访问非顺序，B访问是顺序的，可以说第一种方式读取和写入分别是合并和非合并的，反之亦然。
        且在执行过程中第一种方式(写入非合并)是第二种方式(读取非合并)执行时间的两倍。原因在于读取数据虽然是合并的，但利用了只读数据缓存加载函数__ldg()。从帕斯卡架构开始，若一个编译器能判断一个全局变量在整个核函数中的都只可读(这里的矩阵A),则会自动用函数__ldg()读取全局内存，从而对数据的读取进行缓存，环节非合并访问带来的影响，对于全局写入则没有类似的函数可以利用，一般来说应做到合并的写入。
        对于开普勒和麦克斯韦架构，默认下不会使用__ldg()函数，从而可以将代码改成
B[ny*gridDim.x*blockDim.x+nx]=__ldg(&A[nx*gridDim.y*blockDim.y+ny])；