cudaMallocHost函数详解

在CUDA2.2以下，仅提供cudaMallocHost函数用于分配页锁定内存，与C语言函数malloc分配分页内存相对应。
而从CUDA2.2开始，页锁定内存增加三种新的类型用于主机多线程的portable，用于高效写回write-combined以及零拷贝的mapping，用cudaHostAlloc进行分配，其中采用四个可选参数标志用以指定使用何种特性：
cudaHostAllocDefault: 默认情况等价于cudaMallocHost().
cudaHostAllocPortable: 用于主机多线程数据共享的portable方式
cudaHostAllocMapped: 通过映射方式分配给CUDA实现zero-copy，在设备上使用cudaHostGetDevicePointer()获取指针，使用cudaHostAllocMapped的方式可以实现主机内存到设备内存的映射，设备端不需要额外开辟内存，直接使用cudaHostGetDevicePointer获取主机内存的指针即可。但如果需要频繁的访问该映射指针的话，会大大降低kernel的与运算速度，因为Pin方式下的数据传输速度远远低于gpu的访存带宽，该速度随硬件不同而不同。
cudaHostAllocWriteCombined: 分配写联合的存储空间，主要用于主机到设备的传输或者通过映射页锁定空间CPU写而设备读的情况。但CPU读的效率不高。
　

使用Malloc分配的内存都是Pageable(交换页)的，而另一个模式就是Pinned(Page-locked)，实质是强制让系统在物理内存中完成内存申请和释放的工作，不参与页交换，从而提高系统效率，需要使用cudaHostAlloc和cudaFreeHost(cudaMallocHost的内存也这样释放)来分配和释放。Pageable(交换页)与Pinned(Page-locked)都是“Write-back”，现在X86/X64CPU，会直接在内部使用一个特别的缓冲区，将写入合并，等没满64B（一个cache line)，集中直接写入一次，越过所有的缓存，而读取的时候会直接从内存读取，同样无视各级缓存。这种最大的用途可以用来在CUDA上准备输入数据，因为它在跨PCI-E传输的时候，可能会更快一些（因为不需要询问CPU的cache数据是否在里面）。如果你不需要用它来给GPU准备输入（cudaMemcpyHostToDevice的copy)数据，那么不应该使用此标志，其是用来准备输入数据用的，而基本不会对这段内存再次读取，因为这种内存读取性很低，用这个的好处1、CPU在内存上填写准备数据的时候，可能有跟好的写入特性（因为WC的合并特性）2、写入完毕开始传输，有更好的PCI-e传过去的性能（因为WC的数据肯定不在各级缓存中的特性）。普通的WB内存，在传输的时候需要查询cache中是否有，这会导致额外的性能损失，而WC无这方面的考虑。总之，准备数据用WC，其他情况，例如存放处理完毕放回的数据，不要用这个，应为处理完毕返回的数据在CPU上访问很慢，因为不在缓存中。以及考虑到在很多机器上性能提升微乎其微，但却1、给你带来了额外的需要设定WC标志的代价。2、万一不小心在CPU上当成了读取的数据，会导致变慢。当然，如果某些设备有明显提升写入与PCI-e传输速度，则还是可以用。

Write-through（直写模式）在数据更新时，同时写入缓存Cache和后端存储。此模式的优点是操作简单；缺点是因为数据修改需要同时写入存储，数据写入速度较慢。
Write-back（回写模式）在数据更新时只写入缓存Cache。只在数据被替换出缓存时，被修改的缓存数据才会被写到后端存储。此模式的优点是数据写入速度快，因为不需要写存储；缺点是一旦更新后的数据未被写入存储时出现系统掉电的情况，数据将无法找回。

_________________________________________________________________________________________________________________________

host内存：分为pageable memory 和 pinned memory

pageable memory： 通过操作系统API（malloc（），new（））分配的存储器空间；、

pinned memory：始终存在于物理内存中，不会被分配到低速的虚拟内存中，能够通过DMA加速与设备端进行通信；

                         cudaHostAlloc(), cudaFreeHost()来分配和释放pinned memory；

                         使用pinned memory优点：主机端-设备端的数据传输带宽高；

                                                             某些设备上可以通过zero-copy功能映射到设备地址空间，从GPU直接访问，省掉主存与显存间进行数据拷贝的工作；

                         pinned memory 不可以分配过多：导致操作系统用于分页的物理内存变， 导致系统整体性能下降；通常由哪个cpu线程分配，就只有这个线程才有访问权限；

                         cuda2.3版本中，pinned memory功能扩充：

                                               portable memory：让控制不同GPU的主机端线程操作同一块portable memory，实现cpu线程间通信；使用cudaHostAlloc（）分配页锁定内存时，加上cudaHostAllocPortable标志；

                                                 write-combined Memory：提高从cpu向GPU单向传输数据的速度；不使用cpu的L1，L2 cache对一块pinned memory中的数据进行缓冲，将cache资源留给其他程序使用；在pci-e总线传输期间不会被来自cpu的监视打断；在调用cudaHostAlloc（）时加上cudaHostAllocWriteCombined标志；cpu从这种存储器上读取的速度很低；

                                                mapped memory：两个地址：主机端地址（内存地址），设备端地址（显存地址）。  可以在kernnel程序中直接访问mapped memory中的数据，不必在内存和显存之间进行数据拷贝，即zero-copy功能；在主机端可以由cudaHostAlloc（）函数获得，在设备端指针可以通过cudaHostGetDevicePointer（）获得；通过cudaGetDeviceProperties（）函数返回的canMapHostMemory属性知道设备是否支持mapped memory；在调用cudaHostAlloc（）时加上cudaHostMapped标志，将pinned memory映射到设备地址空间；必须使用同步来保证cpu和GPu对同一块存储器操作的顺序一致性；显存中的一部分可以既是portable memory又是mapped memory；在执行CUDA操作前，先调用cudaSetDeviceFlags（）（加cudaDeviceMapHost标志）进行页锁定内存映射。

 

constant memory：只读地址空间；位于显存，有缓存加速；64Kb；用于存储需要频繁访问的只读参数 ；只读；使用_constant_ 关键字，定义在所有函数之外；两种常数存储器的使用方法：直接在定义时初始化常数存储器；定义一个constant数组，然后使用函数进行赋值；

 

texture memory：只读；不是一块专门的存储器，而是牵涉到显存、两级纹理缓存、纹理拾取单元的纹理流水线；数据常以一维、二维或者三维数组的形式存储在显存中；缓存加速；可以声明大小比常数存储器大得多；适合实现图像树立和查找表；对大量数据的随机访问或非对齐访问有良好的加速效果；在kernel中访问纹理存储器的操作成为纹理拾取（texture fetching）；纹理拾取使用的坐标与数据在显存中的位置可以不同，通过纹理参照系约定二者的映射方式；将显存中的数据与纹理参照系关联的操作，称为将数据与纹理绑定（texture binding）；显存中可以绑定到纹理的数据有：普通线性存储器和cuda数组；存在缓存机制；可以设定滤波模式，寻址模式等；