不同的应用场景推介使用的OpenCL Paths

（一）不同的应用场景推介使用的OpenCL Paths

(1)clEnqueueWriteBuffer()以及clEnqueueReadBuffer()

(2)在pre-pinned host buffer上使用clEnqueueCopyBuffer()

(3)在device buffer上执行 clEnqueueMapBuffer() and
clEnqueueUnmapMemObject()

(4)host直接访问设备zero copy buffer

(5)GPU直接访问host zero copy memory

1. clEnqueueWriteBuffer()以及clEnqueueReadBuffer()

如果应用程序已经通过malloc 或者mmap分配内存，CL_MEM_USE_HOST_PTR是个理想的选择。
有两种使用这种方式的方法：

第一种：
a. pinnedBuffer = clCreateBuffer( CL_MEM_ALLOC_HOST_PTR or CL_MEM_USE_HOST_PTR )

b. deviceBuffer = clCreateBuffer()

c. void *pinnedMemory = clEnqueueMapBuffer( pinnedBuffer )

d. clEnqueueRead/WriteBuffer( deviceBuffer, pinnedMemory )

e. clEnqueueUnmapMemObject( pinnedBuffer, pinnedMemory )

pinning开销在步骤a产生，步骤d没有任何pinning开销。通常应用立即程序执行a,b,c,e步骤，而在步骤d之后，要反复读和修改pinnedMemory中的数据.

第二种：
clEnqueueRead/WriteBuffer 直接在用户的memory buffer中被使用。在copy(host->device)数据前，首先需要pin(lock page)操作，然后才能执行传输操作。这条path大概是peak interconnect bandwidth的2/3。

2. 在pre-pinned host buffer上使用clEnqueueCopyBuffer()

和1类似，clEnqueueCopyBuffer在pre-pinned buffer上以peak interconnect bandwidth执行传输操作：

a. pinnedBuffer = clCreateBuffer( CL_MEM_ALLOC_HOST_PTR or CL_MEM_USE_HOST_PTR )

b. deviceBuffer = clCreateBuffer()
通过:

c. void *memory = clEnqueueMapBuffer( pinnedBuffer )

d. Application writes or modifies memory.

e. clEnqueueUnmapMemObject( pinnedBuffer, memory )

f. clEnqueueCopyBuffer( pinnedBuffer, deviceBuffer )

或者通过:

g. clEnqueueCopyBuffer( deviceBuffer, pinnedBuffer )

h. void *memory = clEnqueueMapBuffer( pinnedBuffer )

i. Application reads memory.

j. clEnqueueUnmapMemObject( pinnedBuffer, memory )

由于pinned memory驻留在host memroy，所以clMap() 以及 clUnmap()调用不会导致数据传输。cpu可以以host memory带宽来操作这些pinned buffer。

3、在device buffer上执行 clEnqueueMapBuffer() and clEnqueueUnmapMemObject()

对于已经通过malloc和mmap分配空间的buffer，传输开销除了interconnect传输外，还要包括一个memcpy过程，该过程把buffer拷贝进mapped device buffer。

a. Data transfer from host to device buffer.

(1). ptr = clEnqueueMapBuffer( .., buf, .., CL_MAP_WRITE, .. )
由于缓冲被映射为write-only,所以没有数据从device传输到host，映射开销比较低。一个指向pinned host buffer的指针被返回。

(2). 应用程序通过memset(ptr)填充host buffer

memcpy ( ptr, srcptr ), fread( ptr ), 或者直接CPU写， 这些操作以host memory全速带宽读写。

(3). clEnqueueUnmapMemObject( .., buf, ptr, .. )

pre-pinned buffer以peak interconnect速度被传输到GPU device。

b. Data transfer from device buffer to host.

(1). ptr = clEnqueueMapBuffer(.., buf, .., CL_MAP_READ, .. )

这个命令启动devcie到host数据传输，数据以peak interconnect bandwidth传输到一个pre-
pinned的临时缓冲中。返回一个指向pinned memory的指针。

(2). 应用程序读、处理数据或者执行 memcpy( dstptr, ptr ), fwrite (ptr), 或者其它类似的函数时候，由
于buffer驻留在host memory中，所以操作以host memory bandwidth执行。
(3). clEnqueueUnmapMemObject( .., buf, ptr, .. )
由于buffer被映射成只读的，没有实际数据传输，所以unmap操作的cost很低。

4. host直接访问device zero copy buffer

这个访问允许数据传输和GPU计算同时执行（overlapped），在一些稀疏(sparse)的写或者更新情况下，比较有用。

a. 一个device上的 zero copy buffer通过下面的命令被创建：

buf = clCreateBuffer ( .., CL_MEM_USE_PERSISTENT_MEM_AMD, .. )

CPU能够通过uncached WC path直接访问该buffer。 通常可以使用双缓冲机制，gpu在处理一个缓存中的数
据，cpu同时在填充另一个缓存中的数据。

b. 传输从host 到 device

(1). ptr = clEnqueueMapBuffer( .., buf, .., CL_MAP_WRITE, .. )
这个操作花销是比较小的，因为zero copy device buffer 直接被映射到host 地址空间的。
(2).应用传输数据通过memset(ptr),memcpy(ptr,srcptr) 或者直接CPU写。CPU直接写入zero copy device buffer通过互连带宽。这依赖于芯片组，带宽和互连带宽是同一个数量级，一般情况下是比峰值要小的。
(3). clEnqueueUnmapMemObject( .., buf, ptr, .. )同map过程一样，这一过程也是低花销的，因为buffer一直驻留在设备中。

c.如果buffer内容必须被读回来，使用

clEnqueueReadBuffer( .., buf, ..) or clEnqueueCopyBuffer( .., buf, zero copy host buffer, .. ).这个
过程会稍慢的，因为host读通过非高速缓存路径。

5 - GPU直接访问host zero copy memory

这个option允许GPU直接读或者写主机memory。一个GPU kernel能读取主机上的数据而不需要传输，并且直接写回到主机 memory。一个理想的用途是直接kernel执行小的I/Os，或者直接将传输延迟集成到kernel的执行时间里。

A. 应用创建一个zero copy host buffer

buf = clCreateBuffer( .., CL_MEM_ALLOC_HOST_PTR, .. )

B. 接下来应用修改或读zero copy host buffer

(1). ptr = clEnqueueMapBuffer( .., buf, .., CL_MAP_READ |
CL_MAP_WRITE, .. )这个操作是较低消耗的，因为他是一个驻留在host memory的buffer的一个映射。
(2). 应用修改数据通过memset(ptr), memcpy,稀疏或密集的读或写。因为应用正修改一个host buffer, 这个操作是以host memory 带宽条件下。
(3). clEnqueueUnmapMemObject( .., buf, ptr, .. )这个操作是耗时很少的，因为buffer是驻留在host memory上。

C. 应用运行clEnqueueNDRangeKernel()

使用此类作为输入或输出，GPU kernel通过互连带宽访问host memory ，并且数据传输变成了kernel 执行的部分。需要的带宽依赖于平台和芯片组，并且和峰值带宽是同一个数量级，对于独立显卡，要注意的是GPU kernel 执行的带宽是低一个数量级的比kernel访问一个位于device上的常规设备buffer。