首先我们了解一些优化时候的术语及其定义：

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)