CUDA学习之全局内存--part2

cuda知识点可真特喵的多。。。。。。。。。。。。。。。。。。。。。

4.3 内存访问模式

大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。

最大限度地利用全局内存带宽是调控核函数性能的基本。

CUDA执行模型的显著特征之一就是指令必须以线程束为单位进行发布和执行。存储操作也是同样。

在执行内存指令时，线程束中的每个线程都提供了一个正在加载或存储的内存地址。在线程束的32个线程中，每个线程都提出了一个包含请求地址的单一内存访问请求，它由一个或多个设备内存传输提供服务。根据线程束中内存地址的分布，内存访问可以被分成不同的模式。

4.3.1 对齐与合并访问

全局内存通过缓存实现加载和存储

全局内存是一个逻辑层面的模型，我们编程的时候有两种模型考虑：
	一种是逻辑层面的，也就是我们在写程序的时候（包括串行程序和并行程序），写的一维（多维）数组，结构体，定义的变量，这些都是在逻辑层面的；
	一种是硬件角度，就是一块DRAM上的电信号，以及最底层内存驱动代码所完成数字信号的处理。

L1表示一级缓存，每个SM都有自己L1，但是L2是所有SM公用的，除了L1缓存外，还有只读缓存和常量缓存 。

读取粒度

核函数运行时需要从全局内存（DRAM）中读取数据，只有两种粒度 ：

• 128字节

• 32字节

   “粒度”，可以理解为最小单位，也就是核函数运行时每次读内存，哪怕是读一个字节的变量，也要读128字节，或者32字节
  

为何是32字节

SM执行的基础是线程束，当一个SM中正在被执行的某个线程需要访问内存，那么，和它同线程束的其他31个线程也要访问内存，这个基础就表示，即使每个线程只访问一个字节，那么在执行的时候，只要有内存请求，至少是32个字节，所以不使用一级缓存的内存加载，一次粒度是32字节而不是更小。

一行一级缓存是128个字节,它映射到设备内存中一个128字节的对齐段。

---

访问方式决定读取粒度

具体是到底是32还是128还是要看访问方式：

• 使用一级缓存：如果启用一级缓存，那么每次从DRAM上加载数据的粒度是128字节
• 不使用一级缓存：如果不适用一级缓存，只是用二级缓存，那么粒度是32字节。

对于CPU来说，一级缓存或者二级缓存是不能被编程的，但是CUDA是支持通过编译指令停用一级缓存的。

一个内存事务（加载和存储都行）：把一次内存请求——也就是从内核函数发起请求，到硬件响应返回数据的过程。

对齐内存访问与合并内存访问

在优化内存的时候，我们要最关注的是以下两个特性

1. 对齐内存访问
   当一个内存事务的首个访问地址是缓存粒度（32或128字节）的偶数倍的时候被称为对齐内存访问。
   非对齐访问就是除上述的其他情况，非对齐的内存访问会造成带宽浪费。

2. 合并内存访问
   当一个线程束中全部的32个线程访问一个连续的内存块时,就会出现合并内存访问。

对齐合并访问的状态是理想化的，也是最高速的访问方式，当线程束从对齐内存地址开始访问一个连续的内存块，那么对齐合并访问出现了。

为了最大化全局内存访问的理想状态，尽量将线程束访问内存组织成对齐合并的方式，这样的效率是最高的。

举例

1. 对齐与合并内存的加载操作
   一个线程束加载数据，使用一级缓存，并且这个事务所请求的所有数据在一个128字节的对齐的地址段上（所有请求的数据在某个首地址是粒度偶数倍的后128个字节里），具体形式如下图，这里请求的数据是连续的，其实可以不连续，但是不要越界就好。
   
   上面蓝色表示全局内存，下面橙色是线程束要的数据，绿色是对齐的地址段。

2. 不在一个对齐的地址段
   
   上图就是典型的一个线程束，数据分散开了，thread0的请求在128之前，后面还有请求在256之后，所以需要三个内存事务，而利用率，也就是从主存取回来的数据被使用到的比例，只有 $128/(128\times 3)$ 的比例。这个比例低会造成带宽的浪费，最极端的表现，就是如果每个线程的请求都在不同的段，也就是一个128字节的事务只有1个字节是有用的，那么利用率只有 $1/128$

内存事务的优化关键：用最少的事务次数满足最多的内存请求。 事务数量和吞吐量的需求随设备的计算能力变化。

4.3.2 全局内存读取

在SM中,数据通过以下3种缓存/缓冲路径进行传输（具体使用何种方式取决于引用
了哪种类型的设备内存）

1. 一级和二级缓存
2. 常量缓存
3. 只读缓存

常规的路径是一级和二级缓存，需要使用常量和只读缓存的需要在代码中显式声明。

但是提高性能，主要还是要取决于访问模式。
全局内存加载操作是否会通过一级缓存取决于两个因素:

1. 设备的计算能力

2. 编译器选项
   编译器禁用一级缓存的选项是：

   -Xptxas -dlcm=cg
   

   编译器启用一级缓存的选项是：

   -Xptxas -dlcm=ca
   

   当一级缓存被禁用的时候，对全局内存的加载请求直接进入二级缓存，如果二级缓存缺失，则由DRAM完成请求。
   每次内存事务可由一个两个或者四个部分执行，每个部分有32个字节，也就是32，64或者128字节一次（注意前面我们讲到是否使用一级缓存决定了读取粒度是128还是32字节，这里增加的64并不在此情况，所以需要注意）。

内存加载访问模式

内存加载可以分为两类：

• 缓存加载 (启用一级缓存)
• 没有缓存的加载 (禁用一级缓存)

内存访问有以下特点：

1. 是否使用缓存：一级缓存是否介入加载过程
2. 对齐与非对齐的：如果访问的第一个地址是32的倍数
3. 合并与非合并，访问连续数据块则是合并的

1. 缓存加载(启用一级缓存)

缓存加载操作经过一级缓存,在粒度为128字节的一级缓存行上由设备内存事务进行传输。

1. 对齐合并的访问，利用率100%
   

2. 对齐的，但是不是连续的，每个线程访问的数据都在一个块内，但是位置是交叉的，利用率100%
   

3. 连续非对齐的，线程束请求一个连续的非对齐的，32个4字节数据，那么会出现，数据横跨两个块，但是没有对齐，当启用一级缓存的时候，就要两个128字节的事务来完成，总线利用率为50%,并且在这两个事务中加载的字节有一半是未使用的。
   

4. 线程束所有线程请求同一个地址，那么肯定落在一个缓存行范围（缓存行的概念没提到过，就是主存上一个可以被一次读到缓存中的一段数据。），那么如果按照请求的是4字节数据来说，使用一级缓存的利用率是 $4/128=3.125\%$

   

5. 最坏的，线程束中线程请求分散于全局内存中的32个4字节地址。
   尽管线程束请求的字节总数仅为128个字节,但地址要占用N个缓存行(0<N≤32)。完成
   一次内存加载操作需要申请N次内存事务。利用率也是 $1/N$
   

    CPU和GPU的一级缓存有显著的差异:
    1. GPU的一级缓存可以通过编译选项等控制，CPU不可以，
    2. CPU一级缓存优化了时间和空间局部性。GPU一级缓存是专为空间局部性而不是为时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。
   

2. 没有缓存的加载 (禁用一级缓存)

当不使用一级缓存的时候，内存事务的粒度变为32字节，更细粒度的好处是提高利用律

1. 对齐合并访问128字节，不用说，还是最理想的情况，使用4个段，利用率 100%
   

2. 对齐不连续访问128字节，都在四个段内，且互不相同，这样的利用率也是 100%
   

3. 连续不对齐，一个段32字节，所以，一个连续的128字节的请求，即使不对齐，最多也不会超过五个段，所以利用率是 4/5=80%
   如果不明白为啥不能超过5个段，请注意前提是连续的，这个时候不可能超过五段
   

4. 所有线程访问一个4字节的数据，那么此时的利用率是 4/32=12.5%
   

5. 最欢的情况，所有目标数据分散在内存的各个角落，由于请求的128个字节最多落在N个32字节的内存分段内而不是N个128个字节的缓存行内,所以相比于缓存加载,即便是最坏的情况也有所改善。
   

3. 非对齐读取示例

为了说明核函数中非对齐访问对性能的影响，我们对第3章中使用的向量加法代码进行修改，去掉所有的内存加载操作，来指定一个偏移量。

注意在下面的核函数中使用了两种索引。新的索引k由给定的偏移量上移，因为偏移量的值可能会导致加载出现非对齐加载。

只有加载数组A和数组B的操作会用到索引k。对数组C的写操作仍使用原来的索引i，以确保写入访问保持对齐。
为保证修改后核函数的正确性,主机代码也应该做出相应的修改

__global__ void sumArraysGPU(float*a,float*b,float*res,int offset,int n)
{
  int i=blockIdx.x*blockDim.x+threadIdx.x;
  int k=i+offset;
  if(k<n)
    res[i]=a[k]+b[k];
}

void sumArrays(float * a,float * b,float * res,int offset,const int size)
{

    for(int i=0,k=offset;k<size;i++,k++)
    {
        res[i]=a[k]+b[k];
    }

}

编译

nvcc -O3 -arch=sm_60 main.cu -o main

使用值为11的偏移量会导致数组A和数组B的内存加载是非对齐的。在这种情况下，运行时间也是最慢的。

通过观察以全局加载效率为指标的结果，可以验证这些非对齐访问就是性能损失的原因:
使用nvprof获取gld_efficiency指标:

对于非对齐读取的情况(偏移量为11)，全局加载效率减半，这意味着请求的全局内存加载吞吐量加倍。

---

禁用一级缓存对全局内存加载性能有何影响

nvcc -O3 -arch=sm_60 -Xptxas -dlcm=cg main.cu -o main

禁用一级缓存的整体性能略低于缓存访问的性能。缓存缺失对非对齐访问的性能影响更大。

• 如果启用缓存，一个非对齐访问可能将数据存到一级缓存，这个一级缓存用于后续的非对齐内存访问。
• 但是，如果禁用一级缓存，那么每一次非对齐请求需要多个内存事务，并且对将来的请求没有作用。
  

对于非对齐情况，禁用一级缓存使加载效率得到了提高，从49.8%提高到了80%。
由于禁用了一级缓存，每个加载请求是在32个字节的粒度上而不是128个字节的粒度上进行处理，因此加载的字节(但未使用的)数量减少了。

禁用一级缓存的整体加载时间并没有减少，但是全局加载效率提高了。
随着设备占用率的提高，禁用一级缓存的加载可帮助提高总线的整体利用率。对于禁用一级缓存的非对齐加载模式来说，未使用的数据传输量可能会显著减少

4. 只读缓存

只读缓存最初是预留给纹理内存加载使用的。对计算能力为3.5及以上的GPU来说，只读缓存也支持使用全局内存加载代替一级缓存。

只读缓存的加载粒度是32个字节。通常,对分散读取来说，这些更细粒度的加载要优于一级缓存。

有两种方法指导内存从只读缓存读取：

• 使用函数 __ldg
• 在间接引用的指针上使用修饰符

代码：
使用内部函数__ldg来通过只读缓存直接对数组进行读取访问

__global__ void copyKernel(float * in,float* out)
{
    int idx=blockDim*blockIdx.x+threadIdx.x;
    out[idx]=__ldg(&in[idx]);
}

注意函数参数，然后就能强制使用只读缓存了。

4.3.3 全局内存写入

写入相对简单很多。

一级缓存不能用在 Fermi 和 Kepler GPU上进行存储操作，发送到设备前，只经过二级缓存，存储操作在32个字节的粒度上执行，内存事物也被分为一段两端或者四段，如果两个地址在一个128字节的段内但不在64字节范围内，则会产生一个四段的事务(也就是说,执行一个四段事务比执行两个一段事务效果更好)

我们将内存写入也参考前面的加载分为下面这些情况：

1. 对齐的，访问一个连续的128字节范围。存储操作使用一个4段事务完成：
   

2. 内存访问是对齐的，但分散在一个192字节的范围内，不连续，使用3个一段事务来搞定
   

3. 对齐的，在一个64字节的范围内，使用一个两段事务完成。
   

除了非对齐情况(偏移量为11)的存储,所有加载和存储的效率都为100%。非对齐写入的存储效率为80%。当偏移量为11且从一个线程束产生一个128个字节的写入请求时,该请求将由一个四段事务和一个一段事务来实现。因此,128个字节用来请求,160个字节用来加载,存储效率为80%。

非对齐写入的情况(偏移量为11)性能最差。

4.3.4 结构体数组与数组结构体

结构体就是基础数据类型组合出来的新的数据类型，这个新的数据类型在内存中表现是：结构中的成员在内存里对齐的依次排开，

1. 数组结构体（AoS） 就是一个数组，每个元素都是一个结构体；
   用代码表示：AoS

   struct A{
   	float x;
   	float y;
   };
   struct A a[N];
   

   它存储的是空间上相邻的数据(例如,x和y),这在CPU上会有良好的缓存局部性。

2. 结构体数组（SoA） 就是结构体中的成员是数组
   用代码表示：SoA

   SoA
   
   struct A{
       int a[N];
       int b[N]
   }a;
   

用AoS模式在GPU上存储示例数据并执行一个只有x字段的应用程序，将导致50%的带宽损失，因为y值在每32个字节段或128个字节缓存行上隐式地被加载。AoS格式也在不需要的y值上浪费了二级缓存空间。

用SoA模式存储数据充分利用了GPU的内存带宽。由于没有相同字段元素的交叉存取，GPU上的SoA布局提供了合并内存访问，并且可以对全局内存实现更高效的利用。

并行编程范式，尤其是SIMD（单指令多数据）对SoA更友好。
CUDA中普遍倾向于SoA，因为这种内存访问可以有效地合并，而被相同内存操作引用的同字段数据元素在存储时是彼此相邻的。

AoS示例

struct naiveStruct{
    float a;
    float b;
};
__global__ void sumArraysGPU(float*a,float*b,struct naiveStruct* res,int n)
{
  //int i=threadIdx.x;
  int i=blockIdx.x*blockDim.x+threadIdx.x;
  if(i<n)
    res[i].a=a[i]+b[i];
}

int nByte=sizeof(float)*nElem;
float *a_h=(float*)malloc(nByte);
float *b_h=(float*)malloc(nByte);

int nByte_struct=sizeof(struct naiveStruct)*nElem;
struct naiveStruct *res_from_gpu_h=(struct naiveStruct*)malloc(nByte_struct);

float *a_d,*b_d;
CHECK(cudaMalloc((float**)&a_d,nByte));
CHECK(cudaMalloc((float**)&b_d,nByte));

struct naiveStruct* res_d;
CHECK(cudaMalloc((struct naiveStruct**)&res_d,nByte_struct));

initialData(a_h,nElem);
initialData(b_h,nElem);

CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice));


sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d,nElem);
cudaDeviceSynchronize();

cudaMemcpy(res_from_gpu_h,res_d,nByte_struct,cudaMemcpyDeviceToHost);

由于仅仅使用数组结构体的一个成员记录结果，因此,请求加载和存储的50%带宽是未使用的。

SoA示例

struct InnerArray
{
    float x[LEN];
    float y[LEN];
};

__global__ void testInnerArray(InnerArray *data, InnerArray * result,
                               const int n)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n)
    {
        float tmpx = data->x[i];
        float tmpy = data->y[i];

        tmpx += 10.f;
        tmpy += 20.f;
        result->x[i] = tmpx;
        result->y[i] = tmpy;
    }
}

100%的效率说明当处理SoA数据布局时,加载或存储内存请求不会重复。每次访问都由一个独立的内存事务来处理。

4.3.5 性能调整

优化设备内存带宽利用率有两个目标：

1. 对齐合并内存访问，以减少带宽的浪费
2. 足够的并发内存操作，以隐藏内存延迟

第三章我们讲过优化指令吞吐量的核函数，实现并发内存访问量最大化是通过以下方式得到的：

1. 增加每个线程中执行独立内存操作的数量
2. 对核函数启动的执行配置进行试验，已充分体现每个SM的并行性

接下来我们就按照这个思路对程序进行优化试验：展开技术和增大并行性。

展开技术

__global__ void sumArraysGPU(float*a,float*b,float*res,int offset,int n)
{
  //int i=threadIdx.x;
  int i=blockIdx.x*blockDim.x*4+threadIdx.x;
  int k=i+offset;
  if(k+3*blockDim.x<n)
  {
      res[i]=a[k]+b[k];
      res[i+blockDim.x]=a[k+blockDim.x]+b[k+blockDim.x];
      res[i+blockDim.x*2]=a[k+blockDim.x*2]+b[k+blockDim.x*2];
      res[i+blockDim.x*3]=a[k+blockDim.x*3]+b[k+blockDim.x*3];
  }

}

展开技术对性能有非常好的影响。对于这样一个I/O密集型的核函数，充分说明内存访问并行有很高的优先级。

两个对齐的测试示例在性能上仍然优于非对齐访问的情况。

展开并不影响执行内存操作的数量，只影响并发执行的数量。

增大并行性

通过调整块的大小来实现并行性调整，也是前面讲过的套路，我们关注的还是内存利用效率
对于展开核函数而言,最佳的线程块大小为每块有256个线程,与之前测试代码中使用的默认的每块有512个线程相比,线程块的数量加倍了。

1. 测试系统使用Fermi GPU,每个SM最多有8个并发线程块，并且每个SM最多有48个并发线程束。

2. 当采用每个线程块有128个线程的方案时，则每个线程块有4个线程束。

3. 因为一个Fermi SM只可以同时放置8个线程块，所以该核函数被限制每个SM上最多有32个线程束。

4. 这可能会导致不能充分利用SM的计算资源,因为没有达到48个线程束的上限。

当非对齐访问被执行时,可以验证线程块大小对性能的影响。以下结果与对齐访问示例产生的结果类似。
这表明,无论访问是否是对齐的,每个SM中相同的硬件资源限制仍会影响核函数的性能。

最大化带宽利用率

影响设备内存操作性能的因素主要有两个:

1. 有效利用设备DRAM和SM片上内存之间的字节移动：为了避免设备内存带宽的浪费,内存访问模式应是对齐和合并的
2. 当前的并发内存操作数：可通过以下两点实现最大化当前存储器操作数。
   1)展开，每个线程产生更多的独立内存访问
   2)修改核函数启动的执行配置来使每个SM有更多的并行性

4.4 核函数可达到的带宽

在分析核函数性能时，需要注意内存延迟，即完成一次独立内存请求的时间;
内存带宽，即SM访问设备内存的速度，它以每单位时间内的字节数进行测量。

本文要做的就是看看这个核函数对应的问题，其极限效率是多少，在理想效率之下，我们来进行优化，我们本文那矩阵转置来进行研究，看看如何把一种看起来没办法优化的内核，重新设计让它达到更好的性能。

4.4.1 内存带宽

多数内核对带宽敏感，也就是说，工人们生产效率特别高，而原料来的很慢，这限制了生产速度。

全局内存中数据的安排方式和线程束的访问方式都对带宽有显著影响。

一般有如下两种带宽

• 理论带宽：
  理论带宽就是硬件设计的绝对最大值，硬件限制了这个最大值为多少，比如对于不使用ECC的Fermi M2090来说，理论峰值 117.6 GB/s
• 有效带宽
  有效带宽是核函数实际达到的带宽，是测量带宽，可以用下面公式计算:
  

注意吞吐量和带宽的区别：

吞吐量是衡量计算核心效率的，用的单位是每秒多少十亿次浮点运算(gflops)，有效吞吐量其不止和有效带宽有关，还和带宽的利用率等因素有关，当然最主要的还是设备的运算核心。

当然，也有内存吞吐量这种说法，这种说法就是单位时间上内存访问的总量，用单位 GB/s 表示，这个值越大表示读取到的数据越多，但是这些数据不一定是有用的。

4.4.2 矩阵转置问题

就是交换矩阵的坐标

所有的数据，结构体也好，数组也好，多维数组也好，所有的数据，在内存硬件层面都是一维排布的，所以我们这里也是使用一维的数组作为输入输出

串行编程

void transformMatrix2D_CPU(float * MatA,float * MatB,int nx,int ny)
{
  for(int j=0;j<ny;j++)
  {
    for(int i=0;i<nx;i++)
    {
      MatB[i*ny+j]=MatA[j*nx+i];
    }
  }
}

通过这个图能得出一个结论，转置操作：

1. 读：原矩阵行进行读取，请求的内存是连续的，可以进行合并访问
2. 写：写到转置矩阵的列中，访问是交叉的

---

交叉访问是使得内存访问变差的罪魁。但是作为矩阵转置本身，这个是无法避免的。但是在这种无法避免的交叉访问前提下，我们怎么能提升效率就变成了一个有趣的课题：

一种是按行读取按列存储：
另一种则是按列读取按行存储：

1. 如果禁用一级缓存加载，那么这两种实现的性能在理论上是相同的。

2. 如果启用一级缓存，那么第二种实现的性能表现会更好。按列读取操作是不合并的(因此带宽将会浪费在未被请求的字节上)，将这些额外的字节存入一级缓存意味着下一个读操作可能会在缓存上执行而不在全局内存上执行。因为写操作不在一级缓存中缓存，所以对按列执行写操作的例子而言，任何缓存都没有意义。

    在Kepler K10、K20和K20x设备中，这两种方法在性能上没有差别，因为一级缓存不用于全局内存访问。
   

1. 为转置核函数设置性能的上限和下限

利用矩阵复制，为矩阵转置寻找性能上下界

假设没有交叉访问，和全是交叉访问的情况，来给出上限和下限：

• 行读取，行存储来复制矩阵(上限)
• 列读取，列存储来复制矩阵(下限)

__global__ void copyRow(float * MatA,float * MatB,int nx,int ny)//行读取，上限
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx=ix+iy*nx;
    if (ix<nx && iy<ny)
    {
      MatB[idx]=MatA[idx];
    }
}
__global__ void copyCol(float * MatA,float * MatB,int nx,int ny)//列读取，下限，全是交叉访问
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx]=MatA[idx];
    }
}

启用一级负载缓存编译代码:

nvcc -O3 -arch=sm_35 -Xptxas -dlcm=ca matrixcpy.cu -o matrixcpy

2. 朴素转置:读取行与读取列

最naive的两种转置方法，不加任何优化

按行加载按列存储

__global__ void transformNaiveRow(float * MatA,float * MatB,int nx,int ny)
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx_row=ix+iy*nx;
    int idx_col=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx_col]=MatA[idx_row];
    }
}

按列加载按行存储（开启一级缓存时，性能优于前者）

__global__ void transformNaiveCol(float * MatA,float * MatB,int nx,int ny)
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx_row=ix+iy*nx;
    int idx_col=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx_row]=MatA[idx_col];
    }
}

使用NaiveCol方法比NaiveRow方法性能表现得更好。如前面所述,导致这种性能提升的一个可能原因是在缓存中执行了交叉读取。即使通过某一方式读入一级缓存中的数据没有都被这次访问使用到，这些数据仍留在缓存中，在以后的访问过程中可能发生缓存命中。

通过缓存交叉读取能够获得最高的加载吞吐量。在缓存读取的情况下，每个内存请求由一个128字节的缓存行来完成。按列读取数据，使得线程束里的每个内存请求都会重复执行32次(因为交叉读取2048个数据元素)，一旦数据预先存储到了一级缓存中，那么许多当前全局内存读取就会有良好的隐藏延迟并取得较高的一级缓存命中率。

对于NaiveCol实现而言，由于合并写入，存储请求从未被重复执行，但是由于交叉读取，多次重复执行了加载请求。这证明了即使是较低的加载效率，一级缓存中的缓存加载也可以限制交叉加载对性能的负面影响。

吞吐量是可以超过带宽的，因为带宽衡量的是从全局内存到SM的速度极限，而吞吐量是SM获得数据的总量除以时间，而这些数据可以来自一级缓存，而不必千里迢迢从主存读取。

3. 展开转置:读取行与读取列

利用展开技术来提高转置内存带宽的利用率。在这个例子中，展开的目的是为每个线程分配更独立的任务，从而最大化当前内存请求。

以下是一个展开因子为4的基于行的实现。

__global__ void transformNaiveRowUnroll(float * MatA,float * MatB,int nx,int ny)
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x*4;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx_row=ix+iy*nx;
    int idx_col=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx_col]=MatA[idx_row];
      MatB[idx_col+ny*1*blockDim.x]=MatA[idx_row+1*blockDim.x];
      MatB[idx_col+ny*2*blockDim.x]=MatA[idx_row+2*blockDim.x];
      MatB[idx_col+ny*3*blockDim.x]=MatA[idx_row+3*blockDim.x];
    }
}

以下是一个展开因子为4的基于列的实现。

__global__ void transformNaiveColUnroll(float * MatA,float * MatB,int nx,int ny)
{
    int ix=threadIdx.x+blockDim.x*blockIdx.x*4;
    int iy=threadIdx.y+blockDim.y*blockIdx.y;
    int idx_row=ix+iy*nx;
    int idx_col=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
        MatB[idx_row]=MatA[idx_col];
        MatB[idx_row+1*blockDim.x]=MatA[idx_col+ny*1*blockDim.x];
        MatB[idx_row+2*blockDim.x]=MatA[idx_col+ny*2*blockDim.x];
        MatB[idx_row+3*blockDim.x]=MatA[idx_col+ny*3*blockDim.x];
    }
}

通过启用一级缓存，“按列加载按行存储” 获得了更好的有效带宽和整体执行时间。

4. 对角转置:读取行与读取列

在DRAM中内存是分区规划的，如果过多的访问同一个区，会产生排队的现象，也就是要等待，为了避免这种情况，我们最好均匀的访问DRAM的某一段，DRAM的分区是每256个字节算一个分区，所以我们最好错开同一个分区的访问，方法就是调整块的ID。

---

当启用一个线程块的网格时，线程块会被分配给SM。编程模型抽象可能用一个一维或二维布局来表示该网格，但是从硬件的角度来看，所有块都是一维的。每个线程块都有其唯一标识符bid，它可以用网格中的线程块按行优先顺序计算得出:

int bid = blockIdx.x + gridDim.x * blockIdx.y

当启用一个核函数时，线程块被分配给SM的顺序由块ID来确定。一旦所有的SM都被完全占用，所有剩余的线程块都保持不变直到当前的执行被完成。一旦一个线程块执行结束，将为该SM分配另一个线程块。由于线程块完成的速度和顺序是不确定的，随着内核进程的执行，起初通过bid相连的活跃线程块会变得不太连续了。

笛卡尔坐标系下的块坐标：

---

所以选用对角坐标
这个对角转置的目的就是使得读取DRAM位置均匀一点，别都集中在一个分区上，方法是打乱线程块，因为连续的线程块可能访问相近的DRAM地址。

使用如下代码代替blockIdx.x，blockIdx.y

int block_y=blockIdx.x;
int block_x=(blockIdx.x+blockIdx.y)%gridDim.x;

__global__ void transformNaiveRowDiagonal(float * MatA,float * MatB,int nx,int ny)
{
    int block_y=blockIdx.x;
    int block_x=(blockIdx.x+blockIdx.y)%gridDim.x;
    
    int ix=threadIdx.x + blockDim.x*block_x;
    int iy=threadIdx.y + blockDim.y*block_y;
    int idx_row=ix + iy*nx;
    int idx_col=ix*ny + iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx_col]=MatA[idx_row];
    }
}
__global__ void transformNaiveColDiagonal(float * MatA,float * MatB,int nx,int ny)
{
    int block_y=blockIdx.x;
    int block_x=(blockIdx.x+blockIdx.y)%gridDim.x;
    
    int ix=threadIdx.x+blockDim.x*block_x;
    int iy=threadIdx.y+blockDim.y*block_y;
    int idx_row=ix+iy*nx;
    int idx_col=ix*ny+iy;
    if (ix<nx && iy<ny)
    {
      MatB[idx_row]=MatA[idx_col];
    }
}

通过使用对角坐标系来修改线程块的执行顺序，这使基于行的核函数性能得到了大大提升。

但是，基于列的核函数在使用笛卡尔块坐标系仍然比使用对角块坐标系表现得更好。

对角核函数的实现可以通过展开块得到更大的提升，但是这种实现不像使用基于笛卡尔坐标系的核函数那样简单直接。

但是这个速度还没有展开的版本快，甚至没有naive的交叉读取速度快，但书上说的是效率有提高，可能是CUDA升级后的原因吧，或者其他原因的影响，但是DRAM分区会出现排队这种现象值得注意。

5. 使用瘦块来增加并行性

增加并行性最简单的方式是调整块的大小。

进一步对使用基于列的NaiveCol核函数的块大小进行试验

目前最佳的块大小为(8,32)，尽管它与大小为(16,16)的块显示了相同的并行性，但这种性能的提升是由“瘦的”块(8,32)带来的

4.5 使用统一内存的矩阵加法

我们经常见到的，在本地分配内存，然后传输到设备，然后在从设备传输回来

而统一内存的基本思路就是减少指向同一个地址的指针，使用统一内存，就没有这些显式的需求了，而是驱动程序帮我们完成。

具体的做法就是:

CHECK(cudaMallocManaged((float**)&a_d,nByte));
CHECK(cudaMallocManaged((float**)&b_d,nByte));
CHECK(cudaMallocManaged((float**)&res_d,nByte));

使用cudaMallocManaged 来分配内存，这种内存在表面上看在设备和主机端都能访问，但是内部过程和我们前面手动copy过来copy过去是一样的，也就是memcopy是本质，而这个只是封装了一下。

 int nxy = nx * ny;
 int nBytes = nxy * sizeof(float);
 // malloc host memory
 float *A, *B, *hostRef, *gpuRef;
 CHECK(cudaMallocManaged((void **)&A, nBytes));
 CHECK(cudaMallocManaged((void **)&B, nBytes));
 CHECK(cudaMallocManaged((void **)&gpuRef,  nBytes);  );
 CHECK(cudaMallocManaged((void **)&hostRef, nBytes););

 // initialize data at host side
 initialData(A, nxy);
 initialData(B, nxy);
 memset(hostRef, 0, nBytes);
 memset(gpuRef, 0, nBytes);

 // add matrix at host side for result checks
 sumMatrixOnHost(A, B, hostRef, nx, ny);
 // invoke kernel at host side
 int dimx = 32;
 int dimy = 32;
 dim3 block(dimx, dimy);
 dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
 sumMatrixGPU<<<grid, block>>>(A, B, gpuRef, nx, ny);

使用统一内存还是手动控制，运行速度差不多。

影响性能差异的最大因素在于CPU数据的初始化时间——使用托管内存耗费的时间更长。

页面故障，

我们分配的这个统一内存地址是个虚拟地址，对应了主机地址和GPU地址，当我们的主机访问这个虚拟地址的时候，会出现一个页面故障，当CPU要访问位于GPU上的托管内存时，统一内存使用CPU页面故障来出发设备到CPU的数据传输，这里的故障不是坏掉了，而是一种通信方式，类似于中断。

故障数和传输数据的大小直接相关。

启用统一内存相关指标。

nvprof --unified-memory-profiling per-process-device ./main

也可以用nvvp查看统一内存的行为

---

虽然统一内存管理给我们写代码带来了方便而且速度也很快，但是实验表明，手动控制还是要优于统一内存管理