OpenCL 维度和工作项

在OpenCL 平台模型中，我们介绍了OpenCL平台模型。但是对于硬件上的两个概念：计算单元、处理单元，并未与软件上的两个概念：工作项、工作组的关系做详细讲解。现在通过一个例子来类比这几者之间的关系。

某个学校高一的年级，这个年级当中会有多个班级，我们假设班级个数为8。高一年级都有计算机课程，会依次去计算机机房里上机，计算机机房里会有电脑，我们假设电脑数为32。每个在机房里的同学根据机房里黑板上老师布置的任务，都在完成属于自己的任务。

对于这样一个场景中的事物与OpenCL中几个概念的类比为：工作项就好比每位同学，工作组就好比一个班级，多个同学组成一个班级，多个工作项也组成一个工作组；机房里的电脑就好比处理单元，机房就好比计算单元。多个类似机房的计算单元构成了一个OpenCL设备。

上述例子就很好地把OpenCL软件层面的工作项和工作组与OpenCL设备硬件层面的处理单元和计算单元联系起来。

在clEnqueueNDRangeKernel（）函数中，参数指定了内核执行工作项的维度、每个维度上的工作项大小和每个维度上每个工作组里的工作项大小。对于每个工作项，需要通过一个专属于自己的ID来执行与其他工作项不同的任务，这个ID在全局工作项中是唯一的，在一个工作组内也是唯一的。在本节中，我们分别从全局工作项ID和工作组中工作项ID来讲述如何获得这个唯一的ID。

维度和工作项

工作项是OpenCL中执行具体内核任务的最小单元，工作项的数量是我们在调用clEnqueueNDRangeKernel（）函数时人为设定的。每个工作项在所有的工作项中都有一个区别于其他工作项的ID，这个ID的起始编号我们也可以人为设定。工作项虽然执行相同代码，但是通过这个ID来确定每个工作项的任务数据，是典型的SIMT模式。

对于一幅灰度图像Img，通过确定x和y方向的位置，我们可以得到该点位置上的灰度值Img（x,y）。如果我们把x和y方向分别用不同的工作项来执行，要处理完这幅灰度图像，我们可以使用二维的工作项。

对于工作项的数目、ID和维度，在OpenCL中都提供了这样的内置函数来查询这几个值。如下表所示。

为了更好地让读者理解上述函数，我们假定一张256×256的灰度图，x方向全局工作项ID需要从3开始；y方向全局工作项ID从5开始，对于如下代码：

主机端代码：

……
//x方向全局工作项ID为252,y方向全局工作项ID为250
const size_t globalSize[2] = {252, 250};
const size_t offset[2] = {3, 5}; //x方向初始偏移3,y方向初始偏移5
err = clEnqueueNDRangeKernel(cmdqueue, kernel, 2, offset,
                                  globalSize, NULL, 0, NULL, NULL);
……

内核代码：

kernel void OffSetTest(_global uchar ＊output)
{
    int index_x = get_global_id(0);        //值为3～255
    int index_y = get_global_id(1);        //值为5～255
    int size_x = get_global_size(0);       //值为252
    int size_y = get_global_size(1);       //值为250
    int offset_x = get_global_offset(0);   //值为3
int offset_y = get_global_offset(1);       //值为5
int dim_size = get_work_dim();             //值为2
}

上述代码中，x方向全局工作项ID的初始偏移量为3，函数clEnqueueNDRangeKernel（）设定x方向全局工作项数目为252，所以x方向全局工作项ID为：3（0+3）～255（252+3）。同理y方向全局工作项ID为：5～255。