内存管理

CUDA编程的目的是给我们的程序加速，尤其是机器学习，人工智能类的计算，CPU不能高效完成，说白了，我们在控制硬件，控制硬件的语言属于底层语言，比如C语言，最头疼的就是管理内存，python，php这些语言有自己的内存管理机制，c语言的内存管理机制——程序员管理。这样的好处是学起来特别困难，但是学会了又会觉得特别爽，因为自由，你可以随意的控制计算机的计算过程。CUDA是C语言的扩展，内存方面基本集成了C语言的方式，由程序员控制CUDA内存，当然，这些内存的物理设备是在GPU上的，而且与CPU内存分配不同，CPU内存分配完就完事了，GPU还涉及到数据传输，主机和设备之间的传输。
接下来我们要了解的是：

为达到最优性能，CUDA提供了在主机端准备设备内存的函数，并且显式地向设备传递数据，显式的从设备取回数据。

这个函数用过很多次了，唯一要注意的是第一个参数，是指针的指针，一般的用法是首先我们生命一个指针变量，然后调用这个函数：

1
2

float * devMem=NULL;
cudaError_t cudaMalloc((float**) devMem, count)

这里是这样的，devMem是一个指针，定义时初始化指向NULL，这样做是安全的，避免出现野指针，cudaMalloc函数要修改devMem的值，所以必须把他的指针传递给函数，如果把devMem当做参数传递，经过函数后，指针的内容还是NULL。
不知道这个解释有没有听明白，通俗的讲，如果一个参数想要在函数中被修改，那么一定要传递他的地址给函数，如果只传递本身，函数是值传递的，不会改变参数的值。
内存分配支持所有的数据类型，什么int，float。。。这些都无所谓，因为他是按照字节分配的，只要是正数字节的变量都能分配，当然我们根本没有半个字节的东西。
函数执行失败返回：cudaErrorMemoryAllocation.
当分配完地址后，可以使用下面函数进行初始化：

1

cudaError_t cudaMemset(void * devPtr,int value,size_t count)

用法和Memset类似，但是注意，这些被我们操作的内存对应的物理内存都在GPU上。
当分配的内存不被使用时，使用下面语句释放程序。

1

cudaError_t cudaFree(void * devPtr)

注意这个参数一定是前面cudaMalloc类的函数（还有其他分配函数）分配到空间，如果输入非法指针参数，会返回 cudaErrorInvalidDevicePointer 错误，如果重复释放一个空间，也会报错。
目前为止，套路基本和C语言一致。但是，设备内存的分配和释放非常影响性能，所以，尽量重复利用！

这个函数我们前面也反复用到，注意这里的参数是指针，而不是指针的指针，第一个参数dst是目标地址，第二个参数src是原始地址，然后是拷贝的内存大小，最后是传输类型，传输类型包括以下几种：

四种方式，都写在字面上来，唯一有点问题的就是有个host 到host，不知道为啥存在，估计很多人跟我想法一样，可能后面有什么高级的用法。
这个例子也不用说了，前面随便找个有数据传输的都有这两步：从主机到设备，然后计算，最后从设备到主机。
代码省略，来张图：

GPU的内存采用的DDR5制式，2011三星才做出来DDR4的主机内存，但是GPU却一直在使用DDR5，这个具体原因我也不清楚，有兴趣的同学自行去查询，但是我们要说的是GPU的内存理论峰值带宽非常高，对于Fermi C2050 有144GB/s，这个值估计现在的GPU应该都超过了，CPU和GPU之间通信要经过PCIe总线，总线的理论峰值要低很多——8GB/s左右，也就是说所，管理不当，算到半路需要从主机读数据，那效率瞬间全挂在PCIe上了。
CUDA编程需要大家减少主机和设备之间的内存传输。

分配count字节的固定内存，这些内存是页面锁定的，可以直接传输到设备的（翻译的原文写的是：设备可访问的，英文原文是：Since the pinned memory can be accessed directly by the device。应该是翻译问题）这样就是的传输带宽变得高很多。
固定的主机内存释放使用：

1

cudaError_t cudaFreeHost(void *ptr)

我们可以测试一下固定内存和分页内存的传输效率，代码如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67

#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"


void sumArrays(float * a,float * b,float * res,const int size)
{
  for(int i=0;i<size;i+=4)
  {
    res[i]=a[i]+b[i];
    res[i+1]=a[i+1]+b[i+1];
    res[i+2]=a[i+2]+b[i+2];
    res[i+3]=a[i+3]+b[i+3];
  }
}
__global__ void sumArraysGPU(float*a,float*b,float*res)
{
  int i=blockIdx.x*blockDim.x+threadIdx.x;
  res[i]=a[i]+b[i];
}
int main(int argc,char **argv)
{
  int dev = 0;
  cudaSetDevice(dev);

  int nElem=1<<14;
  printf("Vector size:%d\n",nElem);
  int nByte=sizeof(float)*nElem;
  float *a_h=(float*)malloc(nByte);
  float *b_h=(float*)malloc(nByte);
  float *res_h=(float*)malloc(nByte);
  float *res_from_gpu_h=(float*)malloc(nByte);
  memset(res_h,0,nByte);
  memset(res_from_gpu_h,0,nByte);

  float *a_d,*b_d,*res_d;
  // pine memory malloc
  CHECK(cudaMallocHost((float**)&a_d,nByte));
  CHECK(cudaMallocHost((float**)&b_d,nByte));
  CHECK(cudaMallocHost((float**)&res_d,nByte));

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

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

  dim3 block(1024);
  dim3 grid(nElem/block.x);
  sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d);
  printf("Execution configuration<<<%d,%d>>>\n",grid.x,block.x);

  CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost));
  sumArrays(a_h,b_h,res_h,nElem);

  checkResult(res_h,res_from_gpu_h,nElem);
  cudaFreeHost(a_d);
  cudaFreeHost(b_d);
  cudaFreeHost(res_d);

  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);

  return 0;
}

注意这个核函数将会被本篇所有程序使用，今天的关键在于主机分配内存部分，所以核函数就选个最简单的。大家看看效率就好。
使用

1

nvprof ./pine_memory

如果提示错误：

1

Error: CUDA profiling error.

可以改用root权限执行，这时候又发现root没有nvprof程序，所以如图一样，用完整路径执行就好，或者添加到你的path里面。
结果如下：

作为对比，我们改写了代码库中第三个里的参数，使用常规内存拷贝方法，得到的时间如下：

这个结果有点尴尬，固定内存的指标显示的是HtoH，也就是主机到主机的内存拷贝，而常规拷贝显示了HtoD。主机到设备但是看memcpy的速度能看出固定内存耗时确实少一些30:42。
同时也能看到cudaHostAlloc和cudaMalloc的时间接近，当数据增大的时候，这个就有区别了，cudaHostAlloc会慢很多。
结论：
固定内存的释放和分配成本比可分页内存要高很多，但是传输速度更快，所以对于大规模数据，固定内存效率更高。
尽量使用流来使内存传输和计算之间同时进行，第六章详细介绍这部分。

前面我们讲，注意线程之间的内存竞争，因为他们可以同时访问同一个内存地址，现在设备和主机可以同时访问同一个设备地址了，所以，我们要注意主机和设备的内存竞争——当使用零拷贝内存的时候。
零拷贝内存是固定内存，不可分页。可以通过以下函数创建零拷贝内存：

1

cudaError_t cudaHostAlloc(void ** pHost,size_t count,unsigned int flags)

最后一个标志参数，可以选择以下值：

有了UVA，可以不用上面的那个获得设备上访问零拷贝内存的函数了：

1

cudaError_t cudaHostGetDevicePointer(void ** pDevice,void * pHost,unsigned flags);

UVA来了以后，此函数基本失业了。
试验，代码：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15

  float *a_host,*b_host,*res_d;
  CHECK(cudaHostAlloc((float**)&a_host,nByte,cudaHostAllocMapped));
  CHECK(cudaHostAlloc((float**)&b_host,nByte,cudaHostAllocMapped));
  CHECK(cudaMalloc((float**)&res_d,nByte));
  res_from_gpu_h=(float*)malloc(nByte);

  initialData(a_host,nElem);
  initialData(b_host,nElem);

  dim3 block(1024);
  dim3 grid(nElem/block.x);
  sumArraysGPU<<<grid,block>>>(a_host,b_host,res_d);
}

UVA代码主要就是差个获取指针，UVA可以直接使用主机端的地址。

这个函数和前面函数结构一致，注意函数名就好了，参数就不解释了，很明显了已经。
CUDA6.0中设备代码不能调用cudaMallocManaged，只能主机调用，所有托管内存必须在主机代码上动态声明，或者全局静态声明
。后面4.5 我们会详细的研究统一内存寻址。