首先我们了解一些优化时候的术语及其定义:
1、deferred allocation(延迟分配),
在第一次使用memory object传输数据时,runtime才对memory object真正分配空间。 这样减少了资源浪费,但第一次使用时要慢一些[一个context多个设备,一个memory object多个location,见前面的blog]。
2.peak interconntect bandwith(峰值内联带宽)
host和device之间通过PCIE总线传输数据,PCIE2.0的上行、下行带宽都是8Gb/s, 对于我们的程序,能达到3Gb/s就不错了,我的笔记本测试只有1.2Gb/s。
3.Pinning(对内存实施pinning操作)
host memory准备向gpu传输时,都要首先进行pinning,就是lock page(禁止交换到外存),pinning操作有一定的性能开销,开销的大小和pinning的host memory大小有关,越大就开销越大。我们可以把host memory分配到pre pinned memory中减少这种开销。
4.WC(write combined operation)
WC是cpu写固定地址时的一个特性,通过把邻接的写操作绑定到一个cacheline,然后发一个写请求,实现了批量写操作。[Gpu内部也有相似的地址合并操作]
5.uncached access
一些内存区域被配置为uncache access,cpu访问比较慢,但是有利于向device memory传输数据,比如前篇日志提到device visible host memory。
6.USWC(无cache的写绑定)
gpu访问uncached的host memory不会产生cache一致性问题,速度会比较快,cpu写因为WC也比较快,相对来说cpu读会变慢。在APU上,这个操作会提供一个快速的cpu写,gpu读的path。
下面看看buffer的分配及使用:
1.normal buffer
用CL_MEM_READ_ONLY/CL_MEM_WRITE_ONLY/CL_MEM_READ_WRITE标志创建的buffer位于device memory中,GPU能够以很高的bandwidth访问这些它,例如对一些高端的显卡,超过100GB/s,host要访问这些内存,只能通过peak interconntect bandwith(PCIE)。
2. zero copy buffer
这种buffer并不做实际的copy工作(除非特殊指定执行copy操作,比如clEnqueueCopyBuffer)。根据创建buffer的type参数,它可能位于host memory也可能位于device memory。
如果device及操作系统支持zero copy,则下面buffer类型可以使用:
• The CL_MEM_ALLOC_HOST_PTR buffer
– zero copy buffer驻留在host。
– host能够以全带宽访问它。
– device通过interconnect bandwidth访问它。
– 这块buffer被分配在prepinned的host memory中。
• The CL_MEM_USE_PERSISTENT_MEM_AMD buffer is
– zero copy buffer 驻留在GPU device中。
– GPU能全带宽访问它。
– host能够以interconnect带宽访问它 (例如streamed写带宽host->device,低的读带宽,因为没有cache利用)。
– 在host和device之间通过interconnect带宽传输数据。
注意:创建buffer的大小是平台dependience的,比如在某个平台上一个buffer不能超过64M,总的buffer不能超过128M等。
zero copy内存在APU上可以得到很好的效果,cpu可以高速的写,gpu能够高速的读,但因为无cache,cpu读会比较慢。
1. buffer = clCreateBuffer(CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY)
2. address = clMapBuffer( buffer )
3. memset( address ) or memcpy( address ) (if possible, using multiple CPU
cores)
4. clEnqueueUnmapMemObject( buffer )
5. clEnqueueNDRangeKernel( buffer )
对于数据量小的传输,zero copy时延(map,unmap等)通常低于相应的DMA引擎时延。
3. prepinned buffer
pinned buffer类型是CL_MEM_ALLOC_HOST_PTR/CL_MEM_USE_HOST_PTR, buffer初始就被创建在prepinned内存中。 EnqueueCopyBuffer以interconnect带宽在host和device之间传输数据(没有pinned和unpinned开销)。
注意:CL_MEM_USE_HOST_PTR能够把已经存在的host buffer转化到pinned memory中去,但是为了保证传输速度,host buffer必须保证256字节对齐。如果只是用来传输数据的话,CL_MEM_USE_HOST_PTR 类型memory对象会一直为prepinned内存,但是它不能作为kernel参数。如果buffer要在kernel中使用的话,runtime会在device创建一个该buffer cache copy,接下来的copy操作不会通过fast path(要保持cache一致性)。
下面的一些函数支持prepinned memory,注意:读取memory可以使用offset:
• clEnqueueRead/WriteBuffer
• clEnqueueRead/WriteImage
• clEnqueueRead/WriteBufferRect (Windows only)