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 )
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.
零复制 zero copy(零复制)是一种特殊形式的内存映射,它允许你将host内存直接映射到设备内存空间上。其实就是设备可以通过直接内存访问(direct memory access,DMA)方式来访问主机的锁页内存。 锁页主机内存 现代操作系统都支持虚拟内存,操作系统实现虚拟内存的主要方法就是通过分页机制。操作系统将内存中暂时不使用的内容换出到外存(硬盘等大容量存储)上,从而腾出空间存
目录1 [GPU硬件架构及运行机制](https://www.cnblogs.com/timlly/p/11471507.html)2 GPU计算基础知识2.1 kernel 核函数2.2 程序层次结构2.3
CUDA
内置变量3 并行
编程
3.1 线程并行3.2 块并行3.3 线程并行与块并行比较3.4 流并行3.5 性能剖析Visual Profiler
1 GPU硬件架构及运行机制
文章写的很全面
2 GPU计算基础知识
CUDA
编程
需要CPU和GPU协同工作,程序中既包含host程序,又包含device程序,他们分别在CPU和GPU上运行,二者间可以进行通信。
host – CPU及其内存
CUDA
程序优化的最终目的,是以最短的时间,在允许的误差范围内完成给定的计算任务。在这里,“最短的时间”是指整个程序的运行时间,更侧重于计算的吞吐量,而不是单个数据的延迟。在开始考虑使用GPU和CPU协同计算之前,应该先粗略地评估使用
CUDA
是否能达到预想的效果,包括以下几个方面:
目前,GPU的单精度计算性能要远远超过双精度计算性能,整数乘法、除法、求模等运算的指令吞吐量也较为有限。在科学计算中,由于需要处理的数据量巨大,往往只有在采用双精度或者四精度时才能获得可靠的结果。目前,采用Tesla架
到底
CUDA
里最传统的内存拷贝
cuda
Malloc,和显式地在主机创建内存
cuda
HostAlloc再开放给GPU访问,以及全局内存寻址
cuda
MallocManaged,这三种方式,哪种在GPU与CPU之间传输数据时,有更高的效率?
为了回答这个问题,于是有了以下对上述三种方式所创建的内存,在主机与设备之间传输速度的比较代码。......
运行时提供的函数允许使用锁页(也称为固定)主机内存(与 malloc() 分配的常规可分页主机内存相反):
cuda
HostAlloc() 和
cuda
FreeHost() 分配和释放锁页主机内存;
cuda
HostRegister() 将 malloc() 分配的内存范围变为锁页内存(有关限制,请参阅参考手册)。
使用页面锁定的主机内存有几个好处:
锁页主机内存和设备内存之间的复制可以与异步并发执行中提到的某些设备的内核执行同时执行。
在某些设备上,锁页主机内存可以映射
在
CUDA
2.2以下,仅提供
cuda
MallocHost函数用于分配页锁定内存,与C语言函数malloc分配分页内存相对应。
而从
CUDA
2.2开始,页锁定内存增加三种新的类型用于主机多线程的portable,用于高效写回write-combined以及零拷贝的mapping,用
cuda
HostAlloc进行分配,其中采用四个可选参数标志用以指定使用何种特性:
cuda
HostAllocDef
平时我们使用的内存都是Pageable(交换页)的,而另一个模式就是Pinned(Page-locked),实质是强制让系统在物理内存中完成内存申请和释放的工作,不参与页交换,从而提高系统效率,需要使用
cuda
HostAlloc和
cuda
FreeHost来分配和释放。
1、带宽更高
2、内核处理和内存拷贝可同时进行
3、可以内存映射(mapped)
一个线程束由32个连续的线程组成,在一个线程束中,所有的线程按照单指令多线程(SIMT)方式执行;即,所有线程都执行相同的指令,每个线程在私有数据上进行操作。
从逻辑角度来看,线程块是线程的集合,它们可以被组织为一维、二维或三维布局。
从硬件角度来看,线程块是一维线程束的集合。在线程块中线程被组织成一维布局,每32个连续线程组成一个线程束。
为了实现CPU与GPU内存的共享,
cuda
采用了零拷贝内存,它值固定内存的一种,当然,也就是实际存储空间实在cpu上。
零拷贝内存的延迟高,在进行频繁的读写操作时尽量少用,否则会大大降低性能。
*创建固定内存映射
* flags:
cuda
HostAllocDefault: make
cuda
HostAlloc same as "
cuda
MallocHost"
一、零拷贝内存相关知识点
之前学习的
CUDA
知识中,主机不能直接访问设备变量,需要通过
cuda
Memcpy
cuda
Memcpy
cuda
Memcpy函数实现主机与设备间数据拷贝,当然设备也不能直接访问主机变量。
这里介绍的零拷贝内存则是个例外,主机和设备都可以访问零拷贝内存。
注意,零拷贝内存相当于从全局内存中分出的一块独立
CUDA
主机端内存分为两种:可分页内存和页锁定内存
可分页内存即通过
操作系统
API(malloc(),new())分配的存储空间;而需要注意的是页锁定内存(pinned memory)。
页锁定内存是由
CUDA
函数cuddaHostAlloc()在主机内存上分配的,页锁定内存的重要属性是主机的
操作系统
不会对这块内存进行分页和交换操作,确保该内存始终驻留在物理内存中。以后总结一下操作系...
如果你想使用
CUDA
进行
编程
,首先需要确保你的电脑上已经安装了 NVIDIA 的
CUDA
工具包。如果没有安装,可以在 NVIDIA 的官网上下载并安装。
其次,你需要使用支持
CUDA
的
编程
语言,比如 C/C++,Fortran 或 Python。你可以使用常见的
编程
环境,如 Visual Studio、Eclipse 或 PyCharm 等,来开发你的
CUDA
程序。
然后,你可以使用
CUDA
的 C/C++ 扩展,在程序中添加特殊的关键字和函数,来调用
CUDA
内核函数和管理 GPU 资源。
最后,你可以使用 NVIDIA 的 nvcc 编译器将你的程序编译成可以在 GPU 上运行的可执行文件。
总的来说,使用
CUDA
进行
编程
需要以下几个步骤:
1. 安装 NVIDIA 的
CUDA
工具包
2. 使用支持
CUDA
的
编程
语言,如 C/C++,编写程序
3. 使用
CUDA
的 C/C++ 扩展,在程序中调用
CUDA
内核函数和管理 GPU 资源
4. 使用 nvcc 编译器将程序编译成可执行文件
如果你还有其他问题,可以继续提出来,我会尽力帮你解决。