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中都提供了这样的内置函数来查询这几个值。如下表所示。

内核NDRange相关函数(1)

为了更好地让读者理解上述函数,我们假定一张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。

Camera课程

Python教程

Java教程

Web教程

数据库教程

图形图像教程

办公软件教程

Linux教程

计算机教程

大数据教程

开发工具教程