3.3 OpenCL执行模型

OpenCL执行模型允许我们建立一个拓扑系统来协调主处理器和其他能够执行OpenCL 内核的设备。为了让内核执行在设备上，还需要对OpenCL上下文进行设置，进而传递执行命令和数据到设备端。

3.3.1 上下文

OpenCL中上下文为了内核的正确执行，进行协调和内存管理。上下文是为了协调主机和设备端的交互，管理设备端可用的内存对象，并持续跟踪在设备上创建出来的程序对象和内核对象。上下文对象可以通过OpenCL APIclCreateContext()进行创建。

  1. cl_context
  2. clCreateContext(
  3.   const cl_context_properties *properties,
  4.   cl_uint num_devices,
  5.   const cl_device_id *devices,
  6.   void (CL_CALL_BACK *pfn_notify)(
  7.     const char *errinfo,
  8.     const void *private_info,
  9.     size_t cb,
  10.     void *user_data),
  11.   void *user_data,
  12.   cl_int *errcode_ret)

properties参数用于限制上下文作用的范围。这个参数可由特定的平台提供，其能够使能与图像的互用性，或使能其他能力。通过限制给定平台的上下文，允许编程者使用多个平台创建的不同的上下文，并且能在一个平台中混用多个供应商提供的设备。另外，创建上下文时必须要使用设备对象，并且编程者可以设置一个用户回调函数，还可以额外传递一个错误码(需要在错误码对象的生命周期内)用于获取API运行的状态。

OpenCL提供了另一个API也能用来创建上下文，其能使用设备列表来创建上下文。通过clCreateContextFromType()可以使用所有的设备类型(CPU、GPU和ALL)创建上下文。创建上下文之后，可以通过使用clGetContextInfo()来查询上下文中设备的数量，以及具体的设备对象。OpenCL中可使用上下文查询指定平台对象和设备对象，通过以上的步骤即可创建上下文，这些步骤也可用于任意OpenCL程序。

3.3.2 命令队列

执行模型是指设备端执行的任务，是基于主机端发送的命令。命令指定的行为包括执行内核，进行数据传递和执行同步。某些设备也能自发一些命令，这种设备以及方式，我们在后面的章节中再进行讨论。

命令队列作为一种通信机制，可以让host发请求到对应的device。当host需要device执行任务的时候，就需要一个命令队列。命令队列需要在每个设备上都进行创建，并且命令队列要在上下文的基础上进行创建。host需要将一条命令提交到对应的命令队列中，因为命令队列不是以分发的形式，而是以指定的形式，所以如果平台上有多个设备时，就需要每个设备上创建一个命令队列。OpenCL中clCreateCommandQueueWithProperties()就是用来创建命令队列，且将命令队列与一个device进行关联。

  1. cl_command_queue
  2. clCreateCommandQueueWithProperties(
  3.   cl_context context,
  4.   cl_device_id device,
  5.   cl_command_queue_properties peoperties,
  6.   cl_int *errcode_ret)

peoperties参数是由一个位域值组成，其可使能命令性能分析功能(CL_QUEUE_PROFILING_ENABLE)，以及/或允许命令乱序执行(CL_QUEUE_OUT_OF_DRDER_EXEC_MODE_ENABLE)。这两个功能将在第5章详细讨论。

对于顺序命令队列(默认创建)，会将命令顺序的推入对应的队列中。乱序命令队列允许OpenCL实现不按入队顺序执行对应的命令，这样的执行方式在某种情况下更为高效。如果使用乱序命令队列，其会根据用户指定的依赖关系，按正确的命令依赖顺序进行执行。

任何以clEnqueue开头的OpenCL API都能向命令队列提交一个命令，并且这些API都需要一个命令队列对象作为输入参数。例如，clEnqueueReadBuffer()将device上的数据传递到host，clEnqueueNDRangeKernel()申请一个内核在对应device执行。如何调用这些API将在后续的章节中进行讨论。

除了向命令队列提交命令的API，OpenCL还包括执行栅栏操作API，这种操作可以用来同步命令队列。clFlush()和clFinish()这两个API都能对命令队列进行栅栏操作。其中，clFinish()的调用将会阻塞host上的执行线程，直到命令队列上的所有命令执行完毕，其功能就是和同步栅栏操作一样。clFlush()将会阻塞host上的执行线程，直到命令队列上的命令都从队列上移出。移出命令队列后的命令，就已经提交到device端，不过不一定完全执行完成。这两个API都需要一个命令队列作为参数。

  1. cl_int clFlush(cl_command_queue command_queue);
  2. cl_int clFinish(cl_command_queue command_queue);

3.3.3 事件

OpenCL API中，用来指定命令之间依赖关系的对象称为事件(event)，所有的clEnqueue开头的API，均有三个共同的参数：事件链表的指针，其指定了当前命令依赖的事件列表，等待列表的长度，以及表示当前命令执行的事件指针，这个指针用于依赖该命令的其他命令。使用事件来指定依赖关系的方式将在第5章介绍。

除了能提供命令依赖顺序，还能通过事件对命令执行的状态随时进行查询。当事件所对应的命令处以执行状态时，其状态就会发生变化。命令状态一共有以下6中：

  • Queued：命令处于命令队列中。
  • Submitted：命令从命令队列中移除，已经提交到设备端执行。
  • Ready：命令已经准备好在设备上执行。
  • Running：命令正在设备上执行。
  • Ended：命令已经在设备上执行完成。
  • Complete：所有命令以及其子命令都执行完成。

子命令与设备端入队有关，我们会在下一节进行讨论。当命令成功的执行完成，事件的状态将会被设置为CL_COMPLETE。如果命令非正常终止，事件的状态将会为一个负数值。这种情况下，有非正常终止的命令队列，以及其他在同一上下文上创建的命令队列，都将不能正常使用或运行。查询事件所使用的API为clGetEventInfo()。

除了记录了不同入队命令间的依赖关系，OpenCL也提供了用于和host同步的APIclWaitForEvents()，该API阻塞host的执行线程，等待指定事件队列上的所有命令执行完毕。

  1. cl_int
  2. clWaitForEvents(
  3.   cl_uint num_events,
  4.   const cl_event *event_list)

3.3.4 设备端入队

目前为止，我们所描述的执行模型是依据“老板-职员”式例来说的，host(老板)向device(职员)发送命令。这样的模型提供了一种简单的主从合作模式。不过，在很多情况下任务的分发并不能静态确定——尤其是算法的下一个阶段要依赖上一个阶段的结果。例如，在组合优化的应用中，查询范围的大小决定着工作组的数量，不过，范围大小只有在上一次迭代的时候才能知道。之前版本的OpenCL，处理这种情况通常是使用一个新的内核对象来执行下一阶段的任务。为了满足这个需求，以及提升性能，OpenCL 2.0为执行模型添加了一项新的特性——设备端入队。

执行中的内核现在可以让另外一个内核进入命令队列中(具体可以看图3.5)。这种情况下，正在执行的内核可以称为“父内核”，刚入队的内核称为“子内核”。虽然，父子内核是以异步的方式执行，但是父内核需要在子内核全部结束后才能结束。我们可通过与父内核关联的事件对象来对执行状态进行查询，当事件对象的状态为CL_COMPLETE时，就代表父内核结束执行。设备端的命令队列是无序命令队列，其具有无序命令队列的所有特性。设备端命令会进入到设备端产生的命令队列中，并且使用事件的方式来存储各个命令间的依赖关系。这些事件对象只有执行在设备端的父内核可见。更多有关设备端入队的内容，将会在第5章进行讨论。