CUDA编程——zero copy

转载 2018年04月17日 14:14:22

零复制

  zero copy(零复制)是一种特殊形式的内存映射,它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问(direct memory access,DMA)方式来访问主机的锁页内存。 
  


锁页主机内存

  现代操作系统都支持虚拟内存,操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存(硬盘等大容量存储)上,从而腾出空间存放将要调入内存的信息。这样,系统好像为用户提供了一个比实际内存大得多的存储器,称为虚拟存储器。 
  锁页就是将内存页面标记为不可被操作系统换出的内存。所以设备驱动程序给这些外设编程时,可以使用页面的物理地址直接访问内存(DMA),从而避免从外存到内存的复制操作。CPU 仍然可以访问上述锁页内存,但是此内存是不能移动或换页到磁盘上的。CUDA 中把锁页内存称为pinned host memory 或者page-locked host memory。


锁页主机内存的优势

  使用锁页内存(page-locked host memory)有一些优势:

  • 锁页内存和GPU内存之间的拷贝可以和内核程序同时执行,也就是异步并发执行。
  • 在一些设备上锁页内存的地址可以从主机地址空间映射到CUDA 地址空间,免去了拷贝开销。
  • 在拥有前线总端的系统上,如果主机内存被分配为锁页内存,主机内存和GPU 内存带宽可以达到更高,如果主机内存被分配为Write-Combining Memory,带宽会进一步提升。

然而锁页主机存储器是稀缺资源,所以锁页内存分配得多的话,分配会失败。另外由于减少了系统可分页的物理存储器数量,分配太多的分页锁定内存会降低系统的整体性能


使用锁页主机内存

  在GPU 上分配的内存默认都是锁页内存,这只是因为GPU 不支持将内存交换到磁盘上。在主机上分配的内存默认都是可分页,如果需要分配锁页内存,则需要使用cudaMallocHost() 或者cudaHostAlloc()。释放时需要使用cudaFreeHost() 释放这一块内存。调用常规的C函数释放,可能会崩溃或者出现一些不常见的错误。也可以通过函数cudaHostRegister() 把可分页内存标记为锁页内存。

__host__ ​cudaError_t cudaMallocHost ( void** ptr, size_t size )

__host__ ​cudaError_t cudaHostAlloc ( void** pHost, size_t size, unsigned int  flags )

__host__ ​cudaError_t cudaFreeHost ( void* ptr )

cudaHostAlloc() 多了一个可选形参flags ,功能更强大。flags 的值可以取如下值。

#define cudaHostAllocDefault 0x00
Default page-locked allocation flag

#define cudaHostAllocMapped 0x02
Map allocation into device space

#define cudaHostAllocPortable 0x01
Pinned memory accessible by all CUDA contexts

#define cudaHostAllocWriteCombined 0x04
Write-combined memory

cudaHostRegister() 函数用于把已经的存在的可分页内存注册为分页锁定的。

__host__ ​cudaError_t cudaHostRegister ( void* ptr, size_t size, unsigned int  flags )
  • 1

flags 是一个可选形参,可以取如下值。

#define cudaHostRegisterDefault 0x00
Default host memory registration flag

#define cudaHostRegisterIoMemory 0x04
Memory-mapped I/O space

#define cudaHostRegisterMapped 0x02
Map registered memory into device space

#define cudaHostRegisterPortable 0x01
Pinned memory accessible by all CUDA contexts

下面分别介绍这些flags 的作用。

Portable Memory

  一块锁页内存可被系统中的所有设备使用(一个系统中有多个CUDA设备时)。 启用这个特性需要在调用cudaHostAlloc() 时使用cudaHostAllocPortable 选项,或者在调用cudaHostRegister() 使用cudaHostRegisterPortable 选项。 
  

Write-Combining Memory

  默认情况下,锁页主机存储是可缓存的。可以在调用cudaHostAlloc() 时传入cudaHostAllocWriteCombined 标签使其被分配为写结合的(Write-Combining Memory)。写结合存储不使用L1 和L2 cache,所以程序的其它部分就有更多的缓存可用。此外,写结合内存通过PCI-E 传输数据时不会被监视(snoop),这能够获得高达40%的传输加速。 从主机读取写结合存储非常慢(因为没有使用L1、L2cache),所以写结合存储应当只用于那些主机只写的存储。 
  

Mapped Memory

  一块锁页内存可以在调用cudaHostAlloc() 分配时传入cudaHostAllocMapped 标签或者在使用cudaHostRegister() 注册时使用cudaHostRegisterMapped 标签,把锁页内存地址映射到设备地址空间。这样,这块存储会有两个地址:一个是从cudaHostAlloc() 或malloc() 返回的在主机内存地址空间上;另一个在设备存储器上,可以通过cudaHostGetDevicePointer() 取得。内核函数可以使用这个指针访问这块存储。 cudaHostAlloc() 返回的地址指针一个的例外情况是,主机和设备使用统一地址空间(Unified Virtual Address Space)。 
内核直接存取主机内存有很多优势:

  • 无需在设备上分配内存,也无需在主机内存和设备内存之间拷贝数据。数据传输是在内核需要的时候隐式进行的。
  • 无须使用流(cuda stream)就可以并发数据传输和内核执行;数据传输和内核执行自动并发执行。

因为映射的锁页主机内存是主机和设备之间共享的,所以在使用cuda stream 或者cuda event 时必须对内存读写同步;避免潜在的写后读,读后写或者写后写等多线程同步问题。 
  为了能够对任何映射的锁页主机内存解引用设备指针,必须在调用任何cuda 运行时函数前调用cudaSetDeviceFlags(),并传入cudaDeviceMapHost 标签。否则,cudaHostGetDevicePointer() 将会返回错误。 
  如果设备不支持被映射分页锁定存储,cudaHostGetDevicePointer() 将会返回错误。程序员可以检查canMapHostMemory 属性,如果设备支持映射锁页主机内存,将会返回1。

注意:使用映射锁页主机内存看,原子操作将不再保证原子性。cudaHostRegisterIoMemory 是cudaHostRegister() 特有的选项,可以把主机内存映射到IO 地址空间。


参考文献

[1]https://en.wikipedia.org/wiki/CUDA_Pinned_memory 
[2] Cook, Shane (2013). CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs (1st ed.). Morgan Kaufmann Publishers Inc. pp. 334–335. ISBN 9780124159334. 

  

原文地址:https://blog.csdn.net/junparadox/article/details/50633641

CUDA零拷贝内存(zerocopy memory)

为了实现CPU与GPU内存的共享,cuda采用了零拷贝内存,它值固定内存的一种,当然,也就是实际存储空间实在cpu上。 零拷贝内存的延迟高,在进行频繁的读写操作时尽量少用,否则会大大降低性能。 /...
  • Rong_Toa
  • Rong_Toa
  • 2017-11-29 16:07:25
  • 382

Zero Copy 零拷贝 简介

 许多web应用都会向用户提供大量的静态内容,这意味着有很多data从硬盘读出之后,会原封不动的通过socket传输给用户。这种操作看起来可能不会怎么消耗CPU,但是实际上它是低效的:kerna...
  • u011591115
  • u011591115
  • 2013-09-26 11:38:25
  • 1937

什么是Zero-Copy?

概述 考虑这样一种常用的情形:你需要将静态内容(类似图片、文件)展示给用户。那么这个情形就意味着你需要先将静态内容从磁盘中拷贝出来放到一个内存buf中,然后将这个buf通过socket传输给用户,进...
  • u013256816
  • u013256816
  • 2016-09-19 21:28:17
  • 7129

Zero-Copy: CUDA, OpenCV and NVidia Jetson TK1

Zero-Copy: CUDA, OpenCV and NVidia Jetson TK1: Part 1 https://ohmwardbond.blogspot.jp/2017/03/zer...
  • u014333051
  • u014333051
  • 2018-01-02 09:28:22
  • 210

CUDA编程模型

参考自《GPU高性能运算之CUDA》主编:张舒。 1、主机与设备        CUDA编程模型将CPU作为主机(Host),GPU作为协处理器(co-processor)或者设备(Device)...
  • sinat_33718563
  • sinat_33718563
  • 2017-06-29 11:25:58
  • 221

CUDA程序优化技巧

有如下几个方面 1. 使用共享内存减少全局内存读取次数; 2. 把全局内存绑定为纹理; 3. 减少bank conflict, 让不同线程读取连续内存,提高cash命中率; 4. 内存对齐,利用GPU...
  • zhangpinghao
  • zhangpinghao
  • 2013-11-18 23:41:35
  • 2335

java-nio之zero copy深入分析

nio零拷贝底层详解
  • u011262847
  • u011262847
  • 2017-09-25 21:04:18
  • 278

理解Netty中的Zero-copy

Zero-copy概念     wiki上关于zero-copy的概念定义     "Zero-copy" describes computer operations in which...
  • qq_26562641
  • qq_26562641
  • 2016-04-26 15:41:32
  • 809

ZMQ zero-copy

零拷贝 第一章中我们曾提过零拷贝是很危险的,其实那是吓唬你的。既然你已经读到这里了,说明你已经具备了足够的知识,能够使用零拷贝。但需要记住,条条大路通地狱,过早地对程序进行优化其实是没有必要的。...
  • bluewind23
  • bluewind23
  • 2014-06-07 11:06:14
  • 617

【基础知识思考整理 】Zero-copy原理理解(用户角度)

基础知识思考整理 http://write.blog.csdn.net/mdeditor#!postId=52836140 关于Zero-Copy的原理。主要参照的是一篇03年的文章[1](L...
  • aganlengzi
  • aganlengzi
  • 2016-11-25 11:53:49
  • 866
收藏助手
不良信息举报
您举报文章:CUDA编程——zero copy
举报原因:
原因补充:

(最多只允许输入30个字)