OpenCL

OpenCL

GPU  OpenCL 

一、 CUDA vs OpenCL

1. 简介

OpenCL: Open Computing Language，开放计算语言。
OpenCL和CUDA是两种异构计算（此异构平台可由CPU，GPU或其他类型的处理器组成。）的编程模型。

  1. CUDA只支持NVIDIA自家的GPU。
  2. OpenCL最早是由Apple提出，后来交给了Khronos这个开放标准组织。OpenCL 1.0 在2008年底正式由Khronos发布，比CUDA晚了整整一年。

2012年移动图形处理器市场份额，imagenation失去苹果后一落千丈，已被别的公司收购：

enter description here
enter description here

2. 操作步骤

CUDA C加速步骤：

  1. 在device (也就是GPU) 上申请内存
  2. 将host (也就是CPU) 上的数据拷贝到device
  3. 执行CUDA kernel function
  4. 将device上的计算结果传回host
  5. 释放device上的内存

OpenCL操作步骤：

  1. 检测申请计算资源
    • 检测platform, clGetPlatformIDs
    • 检测platform对应的device, clGetDeviceInfo
    • 建立context, clCreateContextFromType
    • 建立command queue, clCreateCommandQueue
    • 在context内申请存储空间, clCreateBuffer
  2. 将host (也就是CPU) 上的数据拷贝到device, clCreateBuffer
  3. OpenCL代码编译
    • 读入OpenCL （kernel function） 源代码，创立program 句柄, clCreateProgramWithSource
    • 编译program, clBuildProgram
    • 创立一个 OpenCL kernel 句柄, clCreateKernel
    • 申明设置 kernel 的 参数, clSetKernelArg
    • 设置NDRange
  4. 运行kernel , clEnqueueNDRangeKernel
  5. 将device上的计算结果传回host, clEnqueueReadBuffer
  6. 释放计算资源
    • 释放kernel, clReleaseKernel
    • 释放program, clReleaseProgram
    • 释放device memory, clReleaseMemObject
    • 释放command queue, clReleaseCommandQueue
    • 释放context, clReleaseContext

procedure
procedure

整体架构如下：

enter description here
enter description here

CUDA C语言与OpenCL的定位不同，或者说是使用人群不同。CUDA C是一种高级语言，那些对硬件了解不多的非专业人士也能轻松上手；而OpenCL则是针对硬件的应用程序开发接口，它能给程序员更多对硬件的控制权，相应的上手及开发会比较难一些。

OpenCL, CUDA C
OpenCL, CUDA C

3. 名词比较

Block: 相当于opencl 中的work-group
Thread：相当于opencl 中的work-item
SP: 相当于opencl 中的PE
SM: 相当于opencl 中的CU
warp: 相当于opencl 中的wavefront(简称wave)，基本的调试单位

4. system tradeoff

各种硬件形态的开发效率与执行效率，而opencl在FPGA上作用就是绿色箭头的方向，可以有效提高FPGA开发效率。

system
system

二、常用API

1. clEnqueueNDRangeKernel

clEnqueueNDRangeKernel
clEnqueueNDRangeKernel

参数：

  1. command_queue,
  2. kernel,
  3. work_dim，使用多少维的NDRange，可以设为1, 2, 3, ..., CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS。
  4. global_work_offset(GWO), 每个维度的偏移，如果不设置默认为0
  5. global_work_size(GWS)，每个维度的索引长度，GWS(1) * GWS(2) * ... * GWS(N) 应该大于等于需要处理的数据量
  6. local_work_size(LWS), 每个维度work-group的大小，如果不设置，系统会自己选择一个合适的大小
  7. num_events_in_wait_list: 执行kernel前需要等待的event个数
  8. event_wait_list: 需要等待的event列表
  9. event: 当前这个命令会返回一个event，以供后面的命令进行同步
    返回：

函数返回执行状态。如果成功， 返回CL_SUCCESS

2. clCreateBuffer

clCreateBuffer
clCreateBuffer

  1. context

  2. flags参数共有9种：

    device权限，默认为可读写：
    CL_MEM_READ_WRITE: kernel可读写
    CL_MEM_WRITE_ONLY: kernel 只写
    CL_MEM_READ_ONLY: kernel 只读

    创建方式:
    CL_MEM_USE_HOST_PTR: device端会对host_ptr位置内存进行缓存，如果有多个命令同时使用操作这块内存的行为是未定义的
    CL_MEM_ALLOC_HOST_PTR: 新开辟一段host端可以访问的内存
    CL_MEM_COPY_HOST_PTR: 在devices新开辟一段内存供device使用，并将host上的一段内存内容copy到新内存上

    host权限，默认为可读写:
    CL_MEM_HOST_WRITE_ONLY：host 只写
    CL_MEM_HOST_READ_ONLY： host只读
    CL_MEM_HOST_NO_ACCESS: host没有访问权限

  3. size是buffer的大小

  4. host_ptr只有在CL_MEM_USE_HOST_PTR， CL_MEM_COPY_HOST_PTR时才有效。

一般对于kernel函数的输入参数，使用CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR可以将host memory拷贝到device memory，表示device只读，位置在device上并进行内存复制，host权限为可读写；
对于输出参数，使用CL_MEM_WRITE_ONLY表示device只写，位置在device上，host权限为可读可写。

如果进行host与device之间的内存传递，可以使用clEnqueueReadBuffer读取device上的内存到host上, clEnqueueWriteBuffer可以将host上内存写到device上。

3. clEnqueueWriteBuffer

clEnqueueWriteBuffer
clEnqueueWriteBuffer

  1. command_queue,
  2. buffer, 将内存写到的位置
  3. blocking_write, 是否阻塞
  4. offset, 从buffer的多少偏移处开始写
  5. size, 写入buffer大小
  6. ptr, host端buffer地址
  7. num_events_in_wait_list, 等待事件个数
  8. event_wait_list, 等待事件列表
  9. event, 返回的事件

4. clCreateImage

创建一个ImageBuffer:

clCreateImage
clCreateImage

  1. context
  2. flags, 同clCreateBuffer里的flags
  3. image_format, 图像的属性，包含两个变量: image_channel_order, 指定通道数和形式，通常为RGBA；image_channel_data_type, 定义数据类型， CL_UNORM_INT8表示为unsigned规一化的INT8，CL_UNSIGNED_INT8
    表示 为非规一化的unsigned int8
  4. image_desc, 定义图像的维度大小，
  5. host_ptr, 输入图像地址
  6. errorce_ret, 返回状态

5. clEnqueueWriteImage

clEnqueueWriteImage
clEnqueueWriteImage

  1. command_queue
  2. image, 目标图像
  3. block_writing, 是否阻塞，如果TRUE，则阻塞
  4. origin, 图像的偏移，通常为（0, 0, 0）
  5. region, 图像的区域，(width, height, depth)
  6. input_row_pitch，每行字节数，可能有对齐；如果设为0,则程序根据每个像素的字节数 乘以 width 计算
  7. input_slice_pitch，3D图像的2D slice块，如果是1D或2D图像，这个值必须为0
  8. ptr, host端输入源图像地址
  9. num_events_in_wait_list, 需等待事件个数
  10. evnet_wait_list, 需要等待的事件列表
  11. event, 返回这个命令的事件，用于后续使用

Map buffer

将cl_mem映射到CPU可访问的指针：

clEnqueueMapBuffer
clEnqueueMapBuffer

  1. command_queue
  2. buffer, cl_mem映射的源地址
  3. blocking_map, 是否阻塞
  4. map_flags, CL_MAP_READ，映射的地址为只读；CL_MAP_WRITE，向映射的地址写东西；CL_MAP_WRITE_INVALIDATE_REGION, 向映射的地址为写东西，host不会使用这段地址的内容，这时返回的地址处的内容不保证是最新的
  5. offset, cl_mem的偏移
  6. size, 映射的内存大小
  7. num_events_in_wait_list, 等待事件个数
  8. event_wait_list, 等待事件列表
  9. event, 返回事件
  10. errorcode_ret, 返回状态

返回值是CPU可访问的指针。

注意：

  1. 当flag为CL_MAP_WRITE时，如果不使用unmap进行解映射，device端无法保证可以获取到最新写的值。
  2. 如果不用unmap，那么device端无法释放这部分内存

所以写完内容后，要立马解映射。

buffer

clEnqueueCopyBuffer: 从一个cl buffer拷贝到另一个cl buffer

event

cl_int clWaitForEvents(cl_uint num_events, const cl_event *event_list)
等待事件执行完成才返回，否则会阻塞

cl_int clEnqueueWaitForEvents(cl_command_queue command_queue, cl_uint num_events, const cl_event *event_list)
和 clWaitForEvents 不同的是该命令执行后会立即返回，线程可以在不阻塞的情况下接着执行其它任务。而 clWaitForEvents 会进入阻塞状态，直到事件列表 event_list 中对应的事件处于 CL_COMPLETE 状态。

cl_int clFlush(cl_command_queue command_queue)
只保证command_queue中的command被commit到相应的device上，不保证当clFlush返回时这些command已经执行完。

cl_int clFinish(cl_command_queue command_queue)
clFinish直到之前的队列命令都执行完才返回。clFinish is also a synchronization point.

cl_int clEnqueueBarrier(cl_command_queue command_queue)
屏障命令保证在后面的命令执行之前，它前面提交到命令队列的命令已经执行完成。
和 clFinish 不同的是该命令会异步执行，在 clEnqueueBarrier 返回后，线程可以执行其它任务，例如分配内存、创建内核等。而 clFinish 会阻塞当前线程，直到命令队列为空（所有的内核执行/数据对象操作已完成）。

cl_int clEnqueueMarker(cl_command_queue command_queue, cl_event *event)
将标记命令提交到命令队列 command_queue 中。当标记命令执行后，在它之前提交到命令队列的命令也执行完成。该函数返回一个事件对象 event，在它后面提交到命令队列的命令可以等待该事件。例如，随后的命令可以等待该事件以确保标记之前的命令已经执行完成。如果函数成功执行返回 CL_SUCCESS。

三、架构

1. Platform Model

1个host加上1个或多个device，1个device由多个compute unit组成，1个compute unit又由多个Processing Elemnet组成。

Platform Model
Platform Model

2. Execution Model

执行模型：

一个主机要使得内核运行在设备上，必须要有一个上下文来与设备进行交互。 一个上下文就是一个抽象的容器，管理在设备上的内存对象，跟踪在设备上 创建的程序和内核。

主机程序使用命令队列向设备提交命令，一个设备有一个命令队列，且与上下文 相关。命令队列对在设备上执行的命令进行调度。这些命令在主机程序和设备上 异步执行。执行时，命令间的关系有两种模式：(1)顺序执行，(2)乱序执行。

内核的执行和提交给一个队列的内存命令会生成事件对象，可以用来控制命令的执行、协调宿主机和设备的运行。

有3种命令类型：
• Kernel-enqueue commands: Enqueue a kernel for execution on a device.（执行kernel函数）
• Memory commands: Transfer data between the host and device memory, between memory objects, or map and unmap memory objects from the host address space.（内存传输）
• Synchronization commands: Explicit synchronization points that define order constraints between commands.（同步点）

命令执行经历6个状态：

  1. Queued: 将command放到CommandQueue
  2. Submitted: 将command从CommandQueue提交到Device
  3. Ready: 当所有运行条件满足，放到Device的WorkPool里
  4. Running: 命令开始执行
  5. Ended: 命令执行结束
  6. Complete: command以及其子command都结束执行，并设置相关的事件状态为CL_COMPLETE

Execution Model
Execution Model

Mapping work-items onto an NDRange：

与CUDA里的grid, block, thread类似，OpenCL也有自己的work组织方式NDRange。NDRange是一个N维的索引空间(N为1, 2, 3...)，一个NDRange由三个长度为N的数组定义，与clEnqueueNDRangeKernel几个参数对应：

  1. global_work_size(GWS)，每个维度的索引长度，GWS(1) * GWS(2) * ... * GWS(N) 应该大于等于需要处理的数据量
  2. global_work_offset(GWO), 每个维度的偏移，如果不设置默认为0
  3. local_work_size(LWS), 每个维度work-group的大小，如果不设置，系统会自己选择较好的结果

如下图所示，整个索引空间的大小为，每个work-group大小为，全局偏移为。
对于一个work-item，有两种方式可以索引：

  1. 直接使用global id
  2. 或者使用work-group进行相关计算，设当前group索引为，group里的local id分别为(s_x, s_y)，那么便有

NDRange index space
NDRange index space

3. Memory Model

不同平台的内存模型不一样，为了可移植性，OpenCL定义了一个抽象模型，程序的实现只需要关注抽象模型，而具体的向硬件的映射由驱动来完成。

Memory Model
Memory Model

主要分为host memory和device memory。而device memory 一共有4种内存：
private memory：是每个work-item各自私有
local memory: 在work-group里的work-item共享该内存
global memory: 所有memory可访问
constant memory: 所有memory可访问，只读，host负责初始化

enter description here
enter description here

4. Program Model

OpenCL支持数据并行，任务并行编程，同时支持两种模式的混合。
分散收集（scatter-gather）：数据被分为子集，发送到不同的并行资源中，然后对结果进行组合，也就是数据并行；如两个向量相加，对于每个数据的+操作应该都可以并行完成。
分而治之（divide-and-conquer）：问题被分为子问题，在并行资源中运行，也就是任务并行；比如多CPU系统，每个CPU执行不同的线程。还有一类流水线并行，也属于任务并行。流水线并行，数据从一个任务传送到另外一个任务中，同时前一个任务又处理新的数据，即同一时刻，每个任务都在同时运行。

enter description here
enter description here

并行编程就要考虑到数据的同步与共享问题。

in-order vs out-of-order:
创建命令队列时，如果没有为命令队列设置 CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE 属性，提交到命令队列的命令将按照 in-order 的方式执行。

OpenCL支持两种同步：
同一工作组内(work-group)工作项(work-item)的同步(实现方式barrier)：
reduction的实现中，需要进行数据同步，所谓reduction就是使用多个数据生成一个数据，如tensorflow中的reduce_mean, reduce_sum等。在执行reduce之前，必须保证这些数据已经是有效的，执行过的，

命令队列中处于同一个上下文中的命令的同步(使用clWaitForEvents，clEnqueueMarker, clEnqueueBarrier 或者执行kernel时加入等待事件列表)。

有2种方式同步：
锁（Locks）：在一个资源被访问的时候，禁止其他访问；
栅栏（Barriers）：在一个运行点中进行等待，直到所有运行任务都完成；（典型的BSP编程模型就是这样）

数据共享：
（1）shared memory
当任务要访问同一个数据时，最简单的方法就是共享存储shared memory（很多不同层面与功能的系统都有用到这个方法），大部分多核系统都支持这一模型。shared memory可以用于任务间通信，可以用flag或者互斥锁等方法进行数据保护，它的优缺点：
优点：易于实现，编程人员不用管理数据搬移；
缺点：多个任务访问同一个存储器，控制起来就会比较复杂，降低了互联速度，扩展性也比较不好。
（2）message passing
数据同步的另外一种模型是消息传递模型，可以在同一器件中，或者多个数量的器件中进行并发任务通信，且只在需要同步时才启动。
优点：理论上可以在任意多的设备中运行，扩展性好；
缺点：程序员需要显示地控制通信，开发有一定的难度；发送和接受数据依赖于库方法，因此可移植性差。

Experiment

1. 向量相加

guru_ge@dl:~/opencl/test$ ./cuda_vector_add
SUCCESS
copy input time: 15438.000000
CUDA time: 23.000000
copy output time: 17053.000000
CPU time: 16259.000000
result is right!
guru_ge@dl:~/opencl/test$ ./main
Device: GeForce GTX 1080 Ti
create input buffer time: 7
create output buffer time: 1
write buffer time: 4017
OpenCL time: 639
read buffer time: 30337
CPU time: 16197
result is right!

guru_ge@dl:~/opencl/test$ ./cuda_vector_add
SUCCESS
copy input time: 59825.000000
CUDA time: 36.000000
copy output time: 67750.000000
CPU time: 64550.000000
result is right!
guru_ge@dl:~/opencl/test$ ./main
Device: GeForce GTX 1080 Ti
create input buffer time: 7
create output buffer time: 1
write buffer time: 52640
OpenCL time: 1634
read buffer time: 80206
CPU time: 66502
result is right!
guru_ge@dl:~/opencl/test$

Reference

  1. https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
  2. https://www.cnblogs.com/hlwfirst/p/5003504.html
  3. http://blog.csdn.net/leonwei/article/details/8909897
  4. https://blog.csdn.net/babyfacer/article/details/6863572
  5. https://blog.csdn.net/xbinworld/article/details/45949629
  6. https://blog.csdn.net/Bob_Dong/article/details/70172165?locationNum=11&fps=1
【推广】 免费学中医，健康全家人
原文地址：https://www.cnblogs.com/gr-nick/p/9379352.html