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

理解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 thenumberof 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、参考

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

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

途次客

2018年5月24日

理解OpenCL中的工作组、工作项的索引相关推荐

深入浅出地理解STM32中的定时器工作原理
深入浅出地理解STM32中的定时器工作原理一.如何实现延时 1 纯硬件电路 2 纯软件编程 3 可编程定时/计数器二.可编程定时/计数器有哪些功能? 三.STM32F103 的定时器有哪些 1 基 ...
VS 中 TFS 的工作项查询功能详解
在团队开发中,一般都要使用一些工具来追踪项目开发进度.在 VS 中可以使用 TFS 的工作项来追踪项目开发进度.下面介绍一下使用 TFS 如何进行工作项的查询. 环境: VS 2017 + TFS 2 ...
在Git中，HEAD，工作树和索引之间有什么区别？
有人能告诉我在Git中HEAD,工作树和索引之间的区别吗? 据我所知,它们都是不同分支的名称. 我的假设是否正确? 编辑我找到了这个单个git存储库可以跟踪任意数量的分支,但是您的工作树只与其中一 ...
java类索引_《Java虚拟机原理图解》1.3、class文件中的访问标志、类索引、父类索引、接口索引集合...
讲完了class文件中的常量池,我们就相当于克服了class文件中最麻烦的模块了.现在,我们来看一下class文件中紧接着常量池后面的几个东西:访问标志.类索引.父类索引.接口索引集合. 1. 访问标 ...
java虚拟机标志_《Java虚拟机原理图解》1.3、class文件中的访问标志、类索引、父类索引、接口索引集合...
讲完了class文件中的常量池,我们就相当于克服了class文件中最麻烦的模块了.现在,我们来看一下class文件中紧接着常量池后面的几个东西:访问标志.类索引.父类索引.接口索引集合. 1. 访问标 ...
在MSF中怎么区分易混淆的工作项类型：Bug、风险和问题（我个人的理解）
工作项是 Visual Studio Team Foundation 用于跟踪工作的分配和状态的数据库记录.CMMI® 过程改进 MSF 过程定义了七个工作项来分配和跟踪工作,他们是:任务.更改请求 ...
TFS中的工作项（六）
TFS中的工作项在VS中打开团队资源管理器工作项有功能,产品积压工作项,任务,测试用例,Bug,障碍,它们之间的关系是:一个产品或应用,有一系列功能,功能最大,可以把一个功能分解成多个产品积压工作 ...
有道词典在word中的加载项导致word使用acobat转化PDF时出现word停止工作
一直很奇怪为什么word转PDF怎么一直搞得word生活不能自理,搞了好久都没找到解决办法原来罪魁祸首是有道德在word中的加载项导致的,最好的解决办法就是不让有道德加载项启动. 首先点word的文件 ...
理解vue中less的scoped和/deep/工作原理
理解vue中less的scoped和/deep/工作原理 scoped /deep/ 实战总结 scoped vue项目一般是单页面.多组件,整个项目共用一个css样式表,有时候我们在写组件的过程中 ...

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

理解OpenCL中的工作组、工作项的索引相关推荐

最新文章

热门文章