CUDA02 - 的内存调度与优化

前面一篇（传送门）简单介绍了CUDA的底层架构和一些线程调度方面的问题，但这只是整个CUDA的第一步，下一个问题在于数据的访存：包括数据以何种形式在CPU/GPU之间进行通信、迁移，以及在GPU内部进行存储、访问。

1 global 、shared 、constant、local

通常来讲，待计算的数据都存放在内存或者硬盘（外部存储设备）中，由CPU来进行调度。想要在device上计算、处理数据，就首先需要将数据转移至CUDA，这样的转移操作通常需要经过数据总线实现。这跟叫做PCI-E的总线连接在cpu和gpu之间，负责信号传输，在早期的CUDA版本中程序员必须清楚的知道数据储存在了这根总线的哪一端。当然相比起CUDA的计算速度，总线的访存速度是很慢的。

上一章介绍了CUDA的底层存储结构。在G80中，一个核心计算单元通过访问不同等级的存储设备，来获取计算资源。这些资源有些是属于线程的，有些是属于SM的，还有一些是全局的。下面写一些这些物理结构对应的软件结构，分成了以下几种：

device shared

以__device__ __shared__为关键词声明的变量会被分配至SM上的shared memory，可以由block内的全部线程所共享，生命周期也随着block的结束而结束。（下图位置），为了编程方便，可以直接写为__shared__

device

以__device__ 为关键词声明的变量会被分配至GPU上的global memory，可以由整个GPU上同一grid内的全部线程所共享，生命周期也随着grid的结束而结束。（在G80上，就是这一块）。global memory是整个GPU上与CPU进行数据交互的主要区域。

constant

以__device__ __constant__为关键词声明的变量也会在global memory内分配空间，但是同时会在SM上的 const-cache上进行缓存。因此使用这部分变量的时候，线程会优先在当前sm的cache上进行查找，如果没有命中再考虑查找global memory。这类数据是只读的，因此不会出现不同cache之间的冲突问题。可以缩写为__constant__。

registers

除此之外，也可以选择不加任何关键字，直接像C中那样声明（比如，int num;），这类变量称作Automatic variables(scalar variables) 。这样声明的变量会使用片上的寄存器进行存储，寄存器是由每一个线程所私有的。访存速度最快，同时可用空间也最小。如果用户自己定义的变量大小超过了CUDA提供的物理容量，CUDA会自动将一部分变量分配至global memory 中进行存储，意味着访存速度会降低很多（早期机器中，约等于200个时钟周期）。

CUDA的设计思路是基于SIMT的，一行声明了scalar variable的代码会在每一个线程的私有寄存器上都声明一个对应变量，当线程生命周期结束，对应的变量也会被销毁。当然，如果Automatic variables的对象为一个数组的话，数组会被分配到localmemory之中。注意在这种情况下，即便数据被分配到了local memory，这一变量的生命周期依旧与non-array 一样，并且也只能被当前线程访问。线程结束后，数据的空间也会被回收。这意味着使用 Automatic array variables其实是一种费力不讨好的策略，应尽量避免使用。

寄存器不仅在硬件上有着更快的访存速度，更重要的一点在于访问私有寄存器是可以并行执行的，不会发生访存冲突。

local

以__device__ __local__为关键字声明的变量也在global memory之中存取。尽管其适用范围仅限于thread内部，但是访存速度和使用global 关键字声明的变量速度是基本一样的。之前说到Automatic variables 中有一部分变量有可能会被分配至全局变量；local memory相当于主动将一部分变量分配至 global内。不推荐使用。

2 访存速度优化

上文记录了device内多种等级的存储器结构。下面介绍一些优化CUDA内访存效率的方案，使得用户可以更好的利用这些内存，达到最大的性能。

2.1 latency hide

考虑下面的情景：指令在CUDA上运行，由warp scheduler每次调度一个warp的线程来到片上执行。假设这次计算任务全部都是访存指令；

对于寄存器的访存是没有延迟的在来到片上的1个时钟周期内就能够完成；而如果访存指令的目标是global mem，情况就会不太一样。发出访存指令后，CUDA这一个warp会需要等待很长时间才能得到访存的结果。在前已经写了warp调度的原理：在这个warp等待数据就绪，被挂起的这段时间中，warp scheduler不会将这个warp分配至片上。

问题在于，随着所有warp都执行到了这条指令（假设没有分支），很有可能出现片上所有的warp都在stall状态，等待数据从global mem返回的情况。为了减少这种不必要的阻塞，一个非常直观的思路是：增加warps数量。线程束的数量越多，每个warp的挂起时间占比也就越长，这样就给予了warp更长的时间等候自己请求的数据准备就绪，相当于用执行thread来填补数据的等候时间，保证总有warp在取指。

这里以G80为例子，简单的做一个计算：

假设目标代码是一段IO密集操作，涉及很高密度的global memory存取，共n条访存指令。
每一条针对global memory的访存指令，在cache失效的情况下需要大约400-600个时钟周期。
一个warp内含32个thread，每4个时钟周期切换一个warp

因此，想使用warp片上的执行时间来掩盖访存过程中的阻塞，用户需要至少设计400/4n个warp，才能实现latency hide。

当然上面的情况下一个SM上只有一个Block。为了达到最大的执行效率，一块sm上可能会分得更多的block。

在早期的一些显卡中，latency hide 所起到的作用是巨大的。在使用之前很有必要计算出一个确定的比例，保证cuda可以以最大速度运转。但是随着后续的显卡优化了global memory 的访存时间，这个优势越来约不明显了。

上图：增加 waps数量可以有效的提高计算速度，但是这种提升会达到一个极限。除了memory latency 之外，这个极限值与其他很多属性也有关系，例如运算指令的执行周期及其强度。为了达到相同的运算性能，warps数量会随着计算强度的增长而先增加，后减少。

2.2 register

既然寄存器访存最快，另一个简单的思路是，多使用寄存器。但问题在于寄存器数量是有限的。尽管在代码逻辑上用户可以定义无限多的寄存器，但是实际上能够在寄存器上进行存储的变量并不多。无限制地将变量设为寄存器类型反而会拉低访存效率，因为编译器会自动将多余出来的那部分变量放进global mem。

以GeForce 8800GT为例，每块SM上一共设计了8192个寄存器，每个寄存器都是线程私有的线程之间不能互相访问，意味着用户设计的线程数量 * 每个线程内寄存器的数量应该小于等于这个值。

Nvidia设计8192（共32KB）这个值有着自己的用意。假设每个Block内设计了16 * 16个thread，每个thread又占用了10个寄存器。这样算下来每个block内都需要10 * 256 = 2560个寄存器。同时注意，一个SM上不是只会有一个BLOCK，有可能有多个Block同时运行在一块SM上。假设这个数值为3，得到的寄存器消耗数量为7680个。

在以上条件下，假设每个thread内寄存器消耗加1，总共就会产生2816*3 = 8448 个寄存器，不再能够满足G80的标准。这样一来调度器就不得不减少每块SM上的block数量为2，相当于整体的执行效率降低了1/3。

每个线程分配10个寄存器，这个数量在正常情况下是满足需求的。据说Nvidia做过调查，去查看不同程序所需的寄存器数量，大部分程序都能够满足这个条件。总的来说寄存器就是用来存放临时变量和频繁存期的数据的。

顺带提一下，CUDA中一般对每个block内的线程数量也有着限制，这个限制个人理解和寄存器数量也有很大关系。考虑到一般的程序段，无论如何都需要使用到数个寄存器（如循环变量i等等），如果线程数量太多，反而也会导致计算速度变慢。

2.3 Coalescing

Coalescing 这个说法很形象，直译为聚合，意思就是将一整块数据整体存放。在说明coalescing前，有必要说明下CUDA从global memory中取数的方式。

已经说过，global memory很大，运行速度较慢，每次取数都要花费很长时间。同时，连接到 global memory的带宽很大（64Byte），意味着每次遇到访存指令取数时，同时可以返回64B的数据。假设取数指令为一条简单的load int，余下的60B并不会浪费，而是将这条int数据所在的一整块儿数据同时返回。

利用这个特性来优化访存时间的方法叫Coalescing。具体来讲，这种方法需要数据的使用者操控多个进程，由0号线程访问的首地址单元格必须为64byte的整数倍，后续地址也必须按序访问，不能交叉，不能错位。使用coalescing不需要额外的编写代码，只要满足了条件，编译器会自动进行优化。

如上图所示，thread1 可以选择不使用address 132 处的数据，但不能够错位。由于数据是整个取出的，错位意味着需要再花费相同的时钟周期取下一块内存，不再满足Coalescing的思想。

当然，需要访存的数据并不能保证总是int形。例如，需要线程访问一个结构体，结构体内存储了长方形两条边的长度信息：

struct matrix
{
    float a;
    float b;
}

尽管这次结构体中包含两个float变量，每个线程都需要读取8B的数据，但CUDA仍然可以使用coalescing，只不过同一个warp完成访存需要花费更长的时间。每次的访存大小控制在4Byte、 8Byte、16Byte都可以实现Coalescing。

即便是当访存的需求大小也不满足以上条件时，也有办法可以优化访存。（例如结构体中含有长、宽、高共12Byte的数据）这时有以下两个思路：

第一个思路是想办法凑齐到4、8、16B。在上述例子中只需要再填补一个空的Byte就可以满足16Byte的coalescing，使用 __align(n)__实现内存对齐。

struct __align__(16) cube { 
    float a, b, c; 
};

这样一来，编译器会自动的在访问的时候填补上后面一位，实现coalescing。当然，这样做会带来一些空间上的浪费。

第二个思路是使用数组结构体，来代替结构体数组：

仍旧使用上面的情景。假设待计算的正方体数量只有1024个，那么使用补齐的方法就产生了1024*4B的额外内存开销。如果不想白白浪费掉这些空间，可以尝试改变一下编程思路，不再将一个正方体的长、宽、高连续存放，而是把长单独放在一起、宽单独放在一起：

struct cube{
    float a[256];
    float b[256];
    float c[256];
};

这样一来，每个线程分别访存长、宽、高，三者都满足4KB的coalescing（只不过这种代码写起来多少有点别扭）。

2.4 减少 bank confilct

Coalescing针对global memory的访存进行了优化，而bank conflict 是一个针对shared memory使用的概念。

shared memory 被主动分成了32块，每一块被称作一个bank。例如，地址0x0000 ~ 0x001F 正好分配在 32 个bank上，意味着0x0000、0x0020位于相同的bank内，相同的bank内地址后四位是相同的。

这样的设计主要是为了方便内存的连续读写。shared memory并没有global memory那样大的带宽，但是每个bank的独立性很强，可以做到同时存取。一个warp内有32个线程，一块shared memory内正好有32块bank，这样设计是有用意的。如果代码设计得当，对于一条load指令，同一时刻片上的32个线程可以同时向shared memory内32个bank并行的进行存取，只要一轮即可完成。

可以看到规避bank conflict的触发条件没有那么严格，线程可以错位访问，只要保证不同时访问相同的bank即可。被两个线程同时访问的bank只能串行的处理访存请求，导致等待时间变长。

同一个bank上出现的冲突数量越多，整体的访存速度会越慢。所以bank conflict并不像coalescing 那样all or nothing，每次优化都有效果。

3 数据迁移

上面主要介绍了GPU内部的数据访存和优化，下面的部分主要写一下在GPU和CPU之间，数据是如何交互的。

GPU和CPU都有着自己独立的内存和缓存，由于冯诺依曼结构的限制，数据一般会首先由CPU读入内存，当轮到CUDA完成计算任务的时候，再由CUDA读入device memory；对于分布式系统来说，CUDA device之间也存在着不少的数据交互，这都给数据迁移带来了很高的要求。这种迁移有很多种模式，程序员可能需要主动的控制转移时机，也可以交给系统自动完成迁移，可以异步预取数据，也可以在需要数据的时候阻塞，等待数据迁移；数据甚至可以不进行迁移，每次需要的时候都直接通过总线访问... 总之，在这一部分也存在不少优化空间。

3.1 pcie & nvlink

之前说到，在早期的GPU之中，只有PCIe总线用来连接各个设备。PCIe比起CPU的DRAM间的速度要慢很多；而当系统中存在一台GPU和多台CPU时，这种瓶颈问题会变得更加严重。这部分首先讲一下PCI-e和NVLINK这两条在设备之间进行连接的总线。

上图：两个GPU必须通过PCIe Switch交换数据。

NVLINK在这种情况下诞生了。NVLINK可以理解成一条两端都是插口的总线。只要对应的设备上拥有NVLINK slot，就可以用NVLINK进行连接。相比起使用PCI-e Switch 交换转发数据，NVLINK能够在设备之间直接连接，大大提高了分布式系统的拓扑程度和计算速度。大部分情况下，这种连接设置在gpu之间（甚至可以在两个GPU之间连接多个NVLINK以加强带宽）。当然，只要应的CPU硬件条件允许，在CPU和GPU之间也可以使用（比如IBM power9）。一条NVLINK内提供了好几条全双工的子链路（相比之下，PCI为半双工；PCIe才做到了全双工），根据一些文章的测试结果，NVlink的效率比PCIe要高出两倍以上，这是相当强大的改进了。（详见参考文献）

上图：在峰值带宽和有效带宽上，两者也有着不小的差别。

上图：DGX-1中的拓扑结构。如果没有NVlink，GPU间的数据交互就必须通过pcie Switch。（另：这篇文章中还提到了不同的通讯模式在当前一些主流人工智能和大数据模型下的表现，感兴趣的同学可以去看看。

3.2 相关函数接口

下面先给出各个函数的接口定义，具体的解释在后文展开。

cudaMalloc( void** devPtr, size_t size )

CUDA版的melloc函数，用于在设备上开辟一段空间，存放数组等变量。同C类似的，开辟完空间后需要使用cudaFree函数对空间进行释放。

cudaMallocManaged ( void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal )

使用效果与cudaMalloc 类似，区别在于cudaMallocManaged 分配的空间会使用UM系统自动调度，一般搭配cudaMemPrefetchAsync使用。

cudaMallocManaged 分配旨在供主机或设备代码使用的内存，并且现在仍在享受这种方法的便利之处，即在实现自动内存迁移且简化编程的同时，而无需深入了解 cudaMallocManaged 所分配统一内存 (UM) 实际工作原理的详细信息。nsys profile 提供有关加速应用程序中 UM 管理的详细信息，并在利用这些信息的同时结合对 UM 工作原理的更深入理解，进而为优化加速应用程序创造更多机会。

cudamemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

用于在主机和设备之间同步数据。程序运行到这一步之后会进入阻塞状态，直到同步完成。

主机到设备：cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevice)；设备到主机：cudaMemcpy(h_A,d_A,nBytes,cudaMemcpyDeviceToHost)

cudaMemcpyAsync ( void* ds。t, const void* src, size_t count, cudaMemcpyKind kind, cudaStream_t stream = 0 )

cudamemcpy 的异步版本，参数列表和cudamemcopy也相似。

cudaMemPrefetchAsync ( const void* devPtr, size_t count, int dstDevice, cudaStream_t stream = 0 )：

cudaMemPrefetchAsyn是一个实现数据异步存取的函数接口。将devPtr指针对应的数据复制到dstDevice对应的设备下。与memcopy不同的是，首先这个函数是非阻塞式的，用于异步存取；其次cudaMemPrefetchAsyn只针对使用cudaMallocManaged分配的内存空间，和managed memory。

3.3 cudaMemCopy & cudaMelloc

cudaError_t cudaMalloc (void **devPtr, size_t size );

cudaMalloc 可以简单的理解为cuda中的melloc函数，调用后开辟的空间位于GPU内，第一个参数填入空间首地址，第二个参数填入对应空间的大小。和CPU对应的，在调用cudaMalloc 开辟空间后，还需要调用cudaFree()将对应的空间回收。cuda melloc 还可以为更高维度的数据分配空间。使用cudaMallocPitch，cudaMalloc3D 可以为二维、三维的数据分配空间。

cudaError_t cudaFree ( void * devPtr )
cudaError_t cudaMemcpy (void * dst, const void * src, size_t count， enum kind)

从src向dst拷贝count 数量的内存。kind为一个枚举形，表达数据拷贝的形式，一共有cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, or cudaMemcpyDeviceToDevice四种。

cudaMemCopy 其实可以看成是c语言中memcopy的cuda版本，重点是可以实现不同设备、主机之间的数据交互，是当前最常使用的数据传输方法。

同理，对于更高维度的矩阵，可以使用cudaMemcpy2D/cudaMemcpy3D来拷贝数据。

cudaMemcpyAsync 功能类似，是cudaMemcpy的异步版本。根据传输数据类型的不同，同步和异步函数在表现上有如下差别：

（涉及到一些pinned mem 和 Um的概念会在下文展开）

3.4 UM与内存的页式管理

melloc函数

先看以下一段官方代码：（删掉了一些不太重要的函数声明和异常处理）

#include <stdio.h>

__global__ void addVectorsInto(float *result, float *a, float *b, int N)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;

  for(int i = index; i < N; i += stride)
    result[i] = a[i] + b[i];
}

int main()
{
  const int N = 2<<24;
  size_t size = N * sizeof(float);

  float *a;
  float *b;
  float *c;

  cudaMallocManaged(&a, size);
  cudaMallocManaged(&b, size);
  cudaMallocManaged(&c, size);

  initWith(3, a, N);	//实现代码不再赘述。就是在CPU下位数组赋初值。
  initWith(4, b, N);
  initWith(0, c, N);

  addVectorsInto<<<1, 1>>>(c, a, b, N);
  cudaDeviceSynchronize();

  cudaFree(a);
  cudaFree(b);
  cudaFree(c);
}

可以看到，数据迁移的大体的流程很简单： melloc---> init ---> kernel ---> cudaFree

在这短短的几行代码中有很大的优化空间，我们从Melloc写起。

我们首先考察C语言中Malloc这一函数的表现。使用malloc函数可以为一个占用空间较大的数组其他数据结构分配空间，这样做的好处是可以动态的控制分配空间的大小。在C中已经学过：malloc调用之后，系统只是开辟一段空间，并没有对数据进行赋值，现在需要将这个概念进行一点拓展。

物理上的内存大小是有限的。在一块大小有限的内存上需要运行多个不同的进程，这就导致了用户进程的地址空间很混乱，且充满随机性，不方便程序的编写。这时候操作系统出现了，操作系统为每一个用户进程创建出了一个虚拟的内存空间，这样在编译、取指的时候用户空间可以很好的处理地址问题，不再需要考虑其他进程的影响。这么做的前提是需要操作系统将虚拟空间向真实空间做一个地址映射，这种映射由页表实现，可以简单的理解为一张map，除了映射到的目标地址之外，还需要一些标记位来表示空间的状态信息，等等。

上图：TLB可以简单的理解为页表的一个缓存（但维持一致性的方式与cache不同），设计在cpu中。

当程序请求的数据在物理内存中不存在时，就会发生page fault。之后操作系统会转而向磁盘中查找信息，来替换当前物理页。

说回melloc这一函数。再C语言程序端执行到这一行指令的时候，操作系统会将数据段的最高地址指针向上移动，但这部分地址空间只存在于虚拟空间下，并不存在于物理内存中，页表中更没有对应的页。在第一次访问这一段虚拟空间时（例如melloc后，需要对该区域进行写操作），操作系统会引发缺页中断，这这段时间中建立起相应的映射关系。（顺带一提：引发缺页中断的同时，TLB也会出现相应的 TLB refill 和TLB invalid例外）

在CUDA与CPU的交互过程中，一个很重要的问题就是页错误。

页表在CPU中存放，而物理内存却包括CPU上的内存空间和GPU上的内存空间（甚至有可能是多太GPU设备）。因此，不论是CPU还是GPU的代码段，都有可能面临所需要的数据在页表中缺页的情况。一旦发生了页错误，意味着程序需要进入阻塞状态，等待物理页从一台设备到另一台设备的迁移完成才能继续执行。随着数据量的增加，这一等待过程逐渐成为了程序运行的瓶颈，值得投入精力进行优化。

UVA & UM

首先需要介绍一下最古早的UVA，Unified Virtual Addressing。UVA提供了一个单一寻址系统，出现也比UM早得多。

在UVA中，nvidia统一了不同设备内存设备，将他们全部映射到一个虚拟地址空间下，使得GPU代码空间下的指针可以直接访问内存，不管内存驻留在哪块GPU的device mem、hostmem、还是onchicp shared mem上。也正是在UVA中，使用cudamemcopy可以不需要再声明输入输出参数到底驻留在哪个设备上，很大程度上简化了编程的难度。

同时，UVA也提出了“zero copy memory”的概念。zero copy mem是一种特殊的内存，被pin在了host 的物理内存页上，当device 需要的时候，可以通过PCI-e远程访问，不再需要使用memcopy。“zero copy mem”也可以看作一种在编程效率上的优化，但是可惜并不能对程序性能起到太大的帮助，因为零拷贝并不是无需拷贝，而是一种隐式异步即时拷贝策略，每次访问的时候还是需要走PCIe总线。如果需要频繁的对CPU数据进行读写，可能会收到很大的性能影响。在此之前，从device访问hostmem需要使用cudaHostAlloc（CUDA2.0）， cudaHostGetDevicePointer（CUDA4.0）

总的来说，UVA解决了一部分GPU向CPU访存的问题，但是这条通路是单向的，且效率不高。host无法通过UVA确定驻留在device上的内存地址，因为考虑到数据一致性和结构设计等问题，CUDA VA上由CUDA 映射来的内存是不允许主机访问的。反过来，从host访问device需要等到CUDA 6.0， UM的出现。

UM：Unified Memory是这一部分的重点，在CUDA6.0 之时提出。某种程度上讲，UM统一内存融合了前面两种数据复制方式的优点（零拷贝和显式的数据转移）：GPU可以自由访问整个系统的内存页，并且同时根据需求自由地以更高的带宽进行数据迁移。

UM是一个很优秀的设计，大大简化了编程的复杂度。但是也因为其不可控性，可能会引发页错误，拖慢程序的运行效率。当数据如何在不同计算部件之间进行迁移不重要、很难隔离工作集合时，UM是一个很好的选择.例如，当需要快速写出能够运行的简单代码、拥有大量数据重用、debug或者数据见关联性很强，等等。除了这些情景外，使用UM虽然不会影响程序的正确性，但常常不是问题的最优解。最主要的原因就是UM分配是由程序自动完成的。而假如在需要数据的时候数据并不在对应的设备上（GPU/CPU），往往引发页错误，造成延迟。这一部分在下文cudaMallocManaged 里展开。

UM的进化史。Pascal为另一个比较重要的一个时间节点（CUDA8），为UM的很多设计提供了硬件层面的支持。

cudaMallocManaged for UM

下面说回cudaMallocManaged()这一函数。相比cudaMelloc，cudaMallocManaged 面向的是统一内存。

cudaMallocManaged调用后，程序会自动的分配由主机、设备都可以访问的程序空间。这一空间的真实物理位置由算法决定，对程序员不可见。在早期的K80等设备上（Pascal 之前），调用 cudaMallocManaged() 会在GPU上分配指定数量的空间大小（除非在多GPU环境下（multi-GPU systems with pre-Pascal GPUs），有些GPU不允许peer-to-peer access，这部分内存会被分配至CPU上）。内存分配的同时也在页表上会为相应的虚拟地址段做出标记，初始化相应的页表项，使得操作系统能够得知那部分空间驻留在GPU上。

在这些架构中，内存分配是在调用核函数时完成的。所有需要访问的物理页面都会在进入核函数后预取至运算部件对应的内存中。这些GPU可以面临更少的page fault，但缺点是需要预取大量的物理页（尽管很大一部分可能是没有必要的）。当然，用户可以通过 cudaStreamAttachMemAsync()来使得物理页只迁移至对应的流或核函数中（默认情况下，会将物理页迁移至所有的流/核函数）。

尽管page不会在GPU运行态下出现，但是在CPU运行态下还是会出现pagefault。每当CPU面临对数据的读写操作时，GPU驱动都需要将驻留在GPU上的物理页面迁移至CPU。

在Pascal等后期GPU架构中， cudaMallocManaged() 不会再分配对应的物理空间，只有通过访问或者预取操作（后续会讲到）才会进行填充。这种情况下，除非被CPU/GPU运算单元访问，在内存中是没有对应页表项的。页表中找不到对应的项会引起页错误。此时，程序需要等待物理页在不同的内存之间进行迁移。物理页可以在不同的处理器内存之间进行迁移，程序也会使用一些算法，使得常用的物理页尽可能多的驻留在对应的内存中，减少页错误的频率。但尽管如此，使用这种方法进行内存分配还是会不可避免地拖慢程序运行效率（尤其是当算法存在大量页面迁移时）。

这样，UM的迁移方法就从Bulk migration 变成了On-demand migration。

在后续发布的架构中，陆续为UM加入了一些新功能，使得其可以更好的避免page fault；在VOLTA中CUDA加入了Access Counters；类似与体系结构中的转移预测；GPU driver设计了一个计数器，用来记录最频繁访问到的物理页号，并使其驻留在CUDA内部，减少交换；而NVLINK2 使得CPU可以直接访问甚至缓存一部分GPU内存，把GPUcache保存在自己的L1之中，增加cache命中率。

例如，下方的例子：

#include <iostream>
#include <math.h>
// CUDA kernel to add elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}
 
int main(void)
{
  int N = 1<<20;
  float *x, *y;
 
  // Allocate Unified Memory -- accessible from CPU or GPU
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));
 
  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }
 
  // Launch kernel on 1M elements on the GPU
  int blockSize = 256;
  int numBlocks = (N + blockSize - 1) / blockSize;
  add<<<numBlocks, blockSize>>>(N, x, y);
 
  // Wait for GPU to finish before accessing on host
  cudaDeviceSynchronize();
 
  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;
 
  // Free memory
  cudaFree(x);
  cudaFree(y);
 
  return 0;
}

程序内容其实很简单，就是y[i] += x[i]。注意x和y都是在利用UM进行分配的，在CPU中初始化并在GPU中计算。分别衡量这段代码在K80与P100上的表现：

其中，右图使用了更新的P100，可以看到kernal运行时间明显的增加了。这是因为数据迁移所花费的时间被计算进了核函数的执行时间（在缺页时stall，等待物理页迁移）。在这个测试代码中，数据在CPU中进行初始化，这意味着CUDA需要面对大量的page fault，stall会花费大量的时间。

总的来讲，使用stall - page fault这种方式来代替预取的模式其实节省了更多时间，减少了不必要的内存迁移。但page fault 本身仍是一个程序执行的瓶颈。

page fault及其一些解决方案

Initialize the Data in a Kernel

//之前的代码中数据全部在cpu中进行初始化，而cudamelloc将其空间分配在了GPU上。所以一个简单的办法是干脆在GPU中进行初始化。

想法很简单，但是实际使用起来这一方案可能并不适用。一方面，初始化操作可能不适合并行处理，如果在kernel上执行却只有一个线程工作的话，其效率可能还不如CPU；更大的问题在于很多数据是没办法在GPU上初始化的，例如存储设备大多数情况下只能由CPU访问。谁让人家是核心呢（摊手

Run It Many Times

增加数据的重用次数也可以有效的改善page fault。每次出现页错误意味着物理页的迁移（从cpu到gpu？？？？还是反方向？？）。将核函数运行多次，全部物理页迁移已经在第一次的时候完成了，这样可以有效的拉低平均执行速度。

Prefetching：使用异步访存的函数cudaMemPrefetchAsync来实现。

cudaMemPrefetchAsync()

数据改成 on-demand 转移后，很大一部分页错误会在数据的使用过程中产生。一个很直观的解决方案诞生了：在数据需要被使用之前，由程序员主动将数据迁移至对应的device。这种预取必须是非阻塞式的，否则优化就没有意义。实现这一功能对应的函数就是cudaMemPrefetchAsync。至于预取的时机只能由程序员自己控制。

CUDA 可通过 `cudaMemPrefetchAsync` 函数，轻松将托管内存异步预取到 GPU 设备或 CPU。以下所示为如何使用该函数将数据预取到当前处于活动状态的 GPU 设备，然后再预取到 CPU：

int deviceId;
cudaGetDevice(&deviceId);                                         // The ID of the currently active GPU device.

cudaMemPrefetchAsync(pointerToSomeUMData, size, deviceId);        // Prefetch to GPU device.
cudaMemPrefetchAsync(pointerToSomeUMData, size, cudaCpuDeviceId); // Prefetch to host. `cudaCpuDeviceId` is a
                                                                  // built-in CUDA variable.

在Nvidia的公开课中举了一个使用Prefetch做内存迁移的例子。感兴趣的同学可以看看其他人整理的运行结果。

仍旧使用我之前做加法的例子，只不过这里还包含了一个check函数。由于check函数在CPU上运行，因此数据涉及到 HostToDevice 和 DeviceToHost 两部分，两部分都可以使用数据预取来进行加速。

链接：https://blog.csdn.net/Felaim/article/details/104339103

http://www.greatlakesconsortium.org/events/GPUMulticore/Chapter4-CudaMemoryModel.pdf

https://www2.eecs.berkeley.edu/Pubs/TechRpts/2016/EECS-2016-143.pdf

https://on-demand.gputechconf.com/gtc/2017/presentation/s7285-nikolay-sakharnykh-unified-memory-on-pascal-and-volta.pdf

《深入浅出谈CUDA》

https://developer.nvidia.com/blog/unified-memory-cuda-beginners/

https://docs.nvidia.com/cuda/cuda-memcheck/index.html

https://developer.nvidia.com/blog/maximizing-unified-memory-performance-cuda/

http://info.nvidianews.com/rs/nvidia/images/NVIDIA NVLink High-Speed Interconnect Application Performance Brief.pdf

https://arxiv.org/abs/1903.04611