OpenCL中的工作组、工作项--转

https://blog.csdn.net/zhouxuanyuye/article/details/80445076

 

理解OpenCL中的工作组、工作项的索引

 ==============================================================

目录结构

1、工作组和工作项

2、一维数据的工作组和工作项

3、深度学习中二维图像的池化（Pooling）

4、参考

 ==============================================================

 

关键词：OpenCL   工作组   工作项    深度学习池化Pooling

本文第1节介绍了OpenCL中的工作组和工作项，摘抄自书本《OpenCL编程指南》，因其在并行计算中比较重要，特别作为重点从理论到实践理解。

第2节以核函数为例，该函数的输出是输入数组数据的两倍，同时输出各个工作项对应的全局ID和局部ID。

第3节以深度学习中的池化为例，介绍了2维图像的工作组和工作项ID，以及简单提及3维内容的工作组和工作项ID。

1、工作组和工作项

OpenCL运行时系统会创建一个整数索引空间，索引空间是N维的值网格，N为1、2或3，又称NDRange。

执行内核的各个实例称为工作项（work-item）。工作项在整个索引空间中由一个全局ID标识，就像学校给学生用学号标识。

工作项组织为工作组（work-group）。全局索引为工作组指定了工作组ID（就像学校给班级编号），工作组内又为工作项指定了局部ID（就像班级里又为学生编了号）。

所以，如何通过坐标方式找到工作项。

（1）可以通过工作项的全局索引

（2）先通过工作组索引号，再通过局部索引号

使用(gx,gy)表示工作项的全局ID，全局索引工具大小为（Gx，Gy）。

所以工作项的坐标范围为[(0,Gx-1), (0,Gy-1)]。

用w表示工作组ID，W工作组各个维度的维度大小。

把一个工作组内的索引空间称为局部索引空间。各个维度的大小用字母L表示，局部ID使用小写字母l表示。

因此，大小为（Gx，Gy）的NDRange索引空间划分为Wx*Wy空间上的工作组，工作组的索引号为（wx,wy）。各个工作组大小为Lx*Ly。

Lx = Gx / Wx

Ly = Gy / Wy

通过工作项的全局ID（gx,gy）来定义工作项，或者通过局部ID（lx，ly）和工作组ID（wx,wy）定义：

gx = wx*Lx + lx

gy = wy*Ly + ly

全局索引空间各个维度起始点从0开始，不过在全局索引空间的起始点定义了一个偏移量，用小写字母o表示

gx = wx*Lx + lx + ox

gy = wy*Ly + ly + oy

图1. 2维NDRange中全局ID、局部ID和工作组索引之间关系的一个例子。索引空间的其他参数在图中定义。阴影方块的全局ID为(gx, gy) = (6, 5)，工作组及局部ID分别为(wx,wy) = (1, 1)和(lx, ly) = (2, 1)。

2、一维数据的工作组和工作项

本例中的核函数如下，功能为将输入数据乘以2后输出。输入数据为8*8的数组，在主函数中将设置64个核一起处理，最后打印出输入数据、全局ID号、局部ID号和输出结果，通过以下语句实现：

           

work_dim：工作项的维度，此处为1维。

global_item_size：给每一个核函数编号的全局变量，此处为8*8=64，ID号从0~63。

local_item_size：4，给所有的核函数分组，每四个一组，每组内的ID从0到3。

clEnqueueNDRangeKernel函数的解析可以参考以下网址：

https://blog.csdn.net/gflytu/article/details/7686130

 

            cl_uint work_dim = 1;
            size_t global_item_size;
            size_t local_item_size;
 
            global_item_size = width*height; /*Global number of work items */
            local_item_size = 4; /* Number ofwork items per work group */
                                /*--> global_item_size / local_item_size which indirectly sets the
                                 numberof workgroups*/
 
            /* Execute Data Parallel Kernel */
            ret =clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL,
                                        &global_item_size,&local_item_size,
                                        0, NULL, NULL);

图2. 工作组共有64/4=16个

以上设置共有16个工作组，每个工作组的局部ID有0,1,2,3四个，注意观察控制台打印输出。

核函数：add_vec.cl

__kernel void add_vec(__global int * data_in,
                      __global int *mem_global_id,
                      __global int *mem_local_id,
                      __global int *data_out,
                      int length)
{
            int i,j;
            int global_id;
            int local_id;
 
            global_id = get_global_id(0);
            local_id  = get_local_id(0);
 
            mem_global_id[global_id] =global_id;
            mem_local_id [global_id] = local_id;
 
 
            for(i=0; i<length; i++)
            {
                       data_out[i] =data_in[i]*2;
            }
 
}
主函数：main.cpp

https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/work_id_1dim/main.cpp

 

图3. 一维数组的输入数据、全局ID数组、局部ID数组和输出数据数组

3、深度学习中二维图像的池化（Pooling）

以下例子为二维数据的例子，以深度学习中的池化（Pooling）为例，在深度学习中，池化可以降低数据维度，聚合数据特征，减少计算量等。池化的方法有多种，如平均值池化，最大值池化。可参考https://blog.csdn.net/silence1214/article/details/11809947

本例使用2*2最大值池化，输入数据8*8，输出数据4*4。

使用2维工作组，打印输出各个工作组全局ID和局部ID，如全局ID为（x，y），则输出表示为xy，在核函数中记为10*x+y。

图4. 2*2 max_pooling例子

图5. 二维图像工作组的全局ID和局部ID

以上通过在主函数中的以下语句实现，将工作维度设置为2，设置各个维度的全局大小分别为8和8，通过设置局部块大小来确定工作组，设置为2，所以工作组一共有（8/2）*（8/2）=16个：

            cl_uint work_dim = 2;
            size_t global_item_size[2] = {8, 8};
            size_t local_item_size[2] = {2, 2};
 
            /* Execute Data Parallel Kernel */
            ret =clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL,
                                        global_item_size,local_item_size,
                                        0,NULL, NULL);
如果需要设置三维工作组ID，则以上同理，另外增加两个数组维度分别控制各个组大小。

以下为本例的核函数核主函数，以及实验结果。

核函数：pooling.cl

 

/**********************************************
function:max_pooling, 2*2
2018/05/24
**********************************************/
 
__kernelvoid pooling(__global int * data_in,
                     __global int *mem_global_id,
                     __global int *mem_local_id,
                     __global int *data_out,
                     int width)
{
            int i,j;
            int global_id_x, global_id_y;
            int local_id_x, local_id_y;
 
            global_id_x = get_global_id(0);
            global_id_y = get_global_id(1);
            local_id_x  = get_local_id(0);
            local_id_y  = get_local_id(1);
 
            //(x,y)
            mem_global_id[global_id_y * width +global_id_x] = global_id_x * 10 + global_id_y;
            mem_local_id [global_id_y * width +global_id_x] = local_id_x * 10  +local_id_y;
 
            if((global_id_x % 2 == 0 )&&(global_id_y % 2 == 0))
            {
                       int index1 = global_id_y* width + global_id_x;
                       int index2 = global_id_y* width + global_id_x + width;
                       int tmp1 =max(data_in[index1],data_in[index1 + 1]);
                       int tmp2 =max(data_in[index2],data_in[index2 + 1]);
                       int tmp  = max(tmp1,tmp2);
                       data_out[global_id_y / 2* width / 2 + global_id_x/2] = tmp ;
            }
}
主函数main.cpp

https://github.com/yywyz/OpenCL-Programming-Examples/blob/master/work_id_2dim_pooling/main.cpp

 

图6. 二维数组的输入数据、全局ID数组、局部ID数组和输出池化结果

 

4、参考

https://blog.csdn.net/zhouxuanyuye/article/details/80445076