OpenCl工作组

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

2018年05月24日 22:51:15 途次客 阅读数：489 标签： OpenCL池化Pooling全局ID局部ID 更多

个人分类： 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

1.  
              cl_uint work_dim = 1;
2.  
   size_t global_item_size;
3.  
   size_t local_item_size;
4.  
    
5.  
   global_item_size = width*height; /*Global number of work items */
6.  
   local_item_size = 4; /* Number ofwork items per work group */
7.  
                                   /*--> global_item_size / local_item_size which indirectly sets the
8.  
                                    numberof workgroups*/
9.  
    
10.  
    /* Execute Data Parallel Kernel */
11.  
    ret =clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL,
12.  
                                            &global_item_size,&local_item_size,
13.  
                                          0, NULL, NULL);

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

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

核函数：add_vec.cl

1.  
   __kernel void add_vec(__global int * data_in,
2.  
   __global int *mem_global_id,
3.  
   __global int *mem_local_id,
4.  
   __global int *data_out,
5.  
   int length)
6.  
   {
7.  
   int i,j;
8.  
   int global_id;
9.  
   int local_id;
10.  
     
11.  
    global_id = get_global_id(0);
12.  
    local_id = get_local_id(0);
13.  
     
14.  
    mem_global_id[global_id] =global_id;
15.  
    mem_local_id [global_id] = local_id;
16.  
     
17.  
     
18.  
    for(i=0; i<length; i++)
19.  
    {
20.  
    data_out[i] =data_in[i]*2;
21.  
    }
22.  
     
23.  
    }

主函数：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个：

1.  
   cl_uint work_dim = 2;
2.  
   size_t global_item_size[2] = {8, 8};
3.  
   size_t local_item_size[2] = {2, 2};
4.  
    
5.  
   /* Execute Data Parallel Kernel */
6.  
   ret =clEnqueueNDRangeKernel(command_queue, kernel, work_dim, NULL,
7.  
                global_item_size,local_item_size,
8.  
   0,NULL, NULL);

如果需要设置三维工作组ID，则以上同理，另外增加两个数组维度分别控制各个组大小。

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

核函数：pooling.cl

1.  
   /**********************************************
2.  
   function:max_pooling, 2*2
3.  
   2018/05/24
4.  
   **********************************************/
5.  
    
6.  
   __kernelvoid pooling(__global int * data_in,
7.  
   __global int *mem_global_id,
8.  
   __global int *mem_local_id,
9.  
   __global int *data_out,
10.  
    int width)
11.  
    {
12.  
    int i,j;
13.  
    int global_id_x, global_id_y;
14.  
    int local_id_x, local_id_y;
15.  
     
16.  
    global_id_x = get_global_id(0);
17.  
    global_id_y = get_global_id(1);
18.  
    local_id_x = get_local_id(0);
19.  
    local_id_y = get_local_id(1);
20.  
     
21.  
    //(x,y)
22.  
    mem_global_id[global_id_y * width +global_id_x] = global_id_x * 10 + global_id_y;
23.  
    mem_local_id [global_id_y * width +global_id_x] = local_id_x * 10 +local_id_y;
24.  
     
25.  
    if((global_id_x % 2 == 0 )&&(global_id_y % 2 == 0))
26.  
    {
27.  
    int index1 = global_id_y* width + global_id_x;
28.  
    int index2 = global_id_y* width + global_id_x + width;
29.  
    int tmp1 =max(data_in[index1],data_in[index1 + 1]);
30.  
    int tmp2 =max(data_in[index2],data_in[index2 + 1]);
31.  
    int tmp = max(tmp1,tmp2);
32.  
    data_out[global_id_y / 2* width / 2 + global_id_x/2] = tmp ;
33.  
    }
34.  
    }

主函数main.cpp

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

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

4、参考

第1节摘抄自《OpenCL编程指南》，其他均为个人实践结果。

如有错误，请指出，谢谢阅读。

这个文章确实很好介绍的

转自：https://blog.csdn.net/zhouxuanyuye/article/details/80445076