OpenCL memory object 传输优化

（一）首先了解一些优化时的术语及其定义：

(1)deferred allocation（延迟分配）

(2)peak interconntect bandwith(峰值内联带宽）

(3)Pinning(对内存实施pinning操作）

(4)WC(write combined operation)

(5)uncached access

(6)USWC(无cache的写绑定）

1、deferred allocation（延迟分配）

     在第一次使用memory object传输数据时，runtime才对memory object真正分配空间。 这样减少了资源浪费，但第一次使用时要慢一些。 

 

2、peak interconntect bandwith(峰值内联带宽）

     host和device之间通过PCIE总线传输数据，PCIE3.0的上行、下行带宽都是16Gb/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)Device normal buffer

(2)Zero copy buffer

(3)Prepinned buffer

1.device 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)