opencl学习（四）

由于opencl内核运行时，clEnqueueNDRangeKernel的第5、6个参数global大小和local大小影响计算效率，甚至执行成功情况，想深入了解一下相关的几个参数。参考平台上运行clinfo的输出，最终在CL_DEVICE_MAX_WORK_GROUP_SIZES这个参数上难以解释。查了很多资料，多数只是说明并不合理。从查阅资料中总结出一下内容，在此分享。

freescale的i.MX6q平台clinfo的输出给初学的我带来很多疑问，请有经验的帮忙解惑！

信息查询函数

cl_int clGetDeviceInfo(cl_device_id device,

cl_device_info param_name,

size_t param_value_size,

void * param_value,

size_t *param_value_size_ret )

参数说明

此函数用来查询OpenCL设备信息。首先介绍其参数：
第一个参数device是clGetDeviceIDs的返回值。
第二个参数param_name是一个枚举常量，标识要查询的设备信息，具体有哪些信息稍后详述。
第三个参数param_value_size声明下一个参数param_value所指向的存储空间的字节大小。这个大小要>=查询参数大小。
第四个参数param_value指向要查询参数返回到的存储空间的地址。NULL时表示忽略。
第五个参数返回查询到的参数的实际的大小。设为NULL则忽略。
下面我们来具体介绍一下可以查询的几个常用参数。
CL_DEVICE_TYPE:OpenCL设备类型。目前支持CL_DEVICE_TYPE_CPU,CL_DEVICE_TYPE_GPU, CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_DEFAULT或以上联合。
CL_DEVICE_VENDOR_ID:一个唯一的供应商识别码。一个唯一设备识别码的例子，PCIe ID。
CL_DEVICE_MAX_COMPUTE_UNITS:OpenCL设备上并行计算单元数目。一个work-group只在一个compute unit上执行。该参数最小为1.
CL_DEVICE_MAX_WORK_ITEM_DEMENSIONS:数据并行执行模块用来声明global和localwork-item IDS的最大维度。该参数最小为3。参考clEnqueueNDRangeKernel，该函数第三个参数work_dim是声明global work-item和work-group中的work-items（specify the global work-items and work items in the work-group.）中使用的维度的数目。work-item应该比0大，且<=CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS.
CL_DEVICE_MAX_WORK_ITEM_SIZES:数组类型，指能被在work-group的每一个维度声明的work-item的最大数目。最小值（1,1,1）。
CL_DEVICE_MAX_WORK_GROUP_SIZE:在一个computeunit中执行一个kernel的work-group中work-item的最大数目。最小值为1。Maximum number of work-items in a work-group executing a kernel ona single compute unit, using the data parallelexecution model. (Refer to clEnqueueNDRangeKernel ). The minimum value is1.
2.
cl_int clGetKernelWorkGroupInfo(cl_kernelkernel ,
cl_device_iddevice ,
cl_kernel_work_group_info param_name,
size_t param_value_size,
void* param_value,
size_t*param_value_size_ret )

参数说明：

此函数返回指定到某一device上的kernel对象信息。同样，先来看参数：

第一个参数要查询的kernel对象。

第二个参数指定与kernel绑定的设备列表中的某一个设备。这个列表就是与kernel绑定的context对应的kernel列表。如果列表中只有一个device，此处device参数可以为NULL。

第三个参数指定要查询参数名称，这也是个枚举值。

第三个参数是要指定的要查询的参数返回的字节数，要>=返回值。

第四个参数指返回值指向内存空间的地址，若设为NULL则忽略。

第五个参数返回实际查询到的参数的大小。

下面来说几个重要的可查询的参数：

CL_KERNEL_WORK_GROUP_SIZE:查询在某一指定设备上执行一个kernel可以使用的最大work-group的size。OpenCL实现会使用资源，这就要求kernel确定work-group大小。

This provides a mechanism for the applicationto query the maximum work-group size that can be used to execute a kernel on aspecificdevice given by device. The OpenCL implementationuses the resource requirements of the kernel (register usage etc.) to determinewhat this work -group size should be.

work-group/work-item/size等关系说明

为执行一个数据并行kernel，除work-items的数目外也要指定work-groups的数目。这也就是为什么两个参数都必须传递给clEnqueueNDRangeKernel。例如：

size_t global_item_size = 4;//总的线程数

size_t local_item_size = 1;//每一个group的线程数

/* Execute OpenCL kernel as data parallel*/

ret = clEnqueueNDRangeKernel(command_queue,kernel, 1, NULL,

&global_item_size,&local_item_size, 0, NULL, NULL);

这个就表示上面这个数据并行计算的kernel中每一个work-group由1个work-item组成，而共有4个work-items要被处理，即总的work-items要被分成4个work-group。

另外work-item对应硬件上的一个PE（processing element）,而一个work-group对应硬件上的一个CU（computing unit）。这种对应可以理解为，一个work-item不能被拆分到多个PE上处理；同样，一个work-group也不能拆分到多个CU上同时处理（忘了哪里看到的信息）。当映射到OpenCL硬件模型上时，每一个work-item运行在一个被称为处理基元（processing element）的抽象硬件单元上，其中每个处理基元可以处理多个work-item(注：摘自《OpenCL异构计算》P87)。（如此而言，是不是说对于二维的globalx必须是localx的整数倍，globaly必须是localy的整数倍？那么如果我数据很大，work-item所能数量很多，如果一个group中中work-item的数量不超过CU中PE的个数，那么group的数量就可能很多；如果我想让group数量小点，那work-item的数目就会很多，还能不能处理了呢？这里总是找不多一个权威的解释，还请高手指点！针对group和item的问题）。

每一个work-group中work-item的数目是不能改变的，始终如一。如果work-item的数目不能在work-groups中均分,clEnqueueNDRangeKernel失败，返回错误码CL_INVALID_WORK_GROUP_SIZE。此处要注意，自己在尝试检测GPU处理能力的时候给出的work-item和work-group的数目不能整除时不一定是数量超限，有可能只是不能整除。

global work-item ID、localwork-item ID，和work-group ID之间的关系如下图所示。

图1 work-group ID和work-item ID

表1 获取ID的函数

函数	返回值
get_group_id	Work-group ID
get_global_id	Global work-item ID
get_local_id	Local work-item ID

因为要处理2D图像或3D空间，work-items和work-groups可以被指定为2或3维。图2给出一个work-group和work-item被定义为2D的例子。

图2 work-group和work-item定义为2D

因为work-group和work-item可至3维，get_group_id(), get_global_id(), get_local_id()每一个的参数可以是0~2。

注意，空间维度指数和每个work-group中work-item的数目能够依据设备而变化。最大维度指数可以通过clGetDeviceInfo()来获取CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS，每个work-group中work-items的最大值可以通过CL_DEVICE_MAX_WORK_ITEM_SIZES获取。前者是cl_uint型，后者是size_t的数组。

在我使用的freescale i.MX6q平台上clEnqueueNDRangeKernel的第6个参数local_item_size不能大于clinfo中输出的CL_KERNEL_WORK_GROUP_SIZE的大小192。最终还是没有解决一个问题：CL_DEVICE_MAX_WORK_GROUP_SIZE（clGetDeivceInfo获取）到底是干嘛用的？在我的平台上该值为1024，它跟CL_KERNEL_WORK_GROUP_SIZE（clGetKernelWorkGroupInfo获取192）什么区别？

注：本人处在学习阶段，博文是自己学习过程中的所得及理解。对初学者不可避免会出现一些错误理解，还请大家批评指正，相互学习。

opencl学习（四）相关推荐

OpenCL学习笔记——整体流程（向量相加）
OpenCL学习笔记--整体流程 OpenCL可以实现混合设备的并行计算,这些设备包括CPU,GPU,以及其他处理器,比如Cell处理器,DSP等.使用OpenCL编程,可以实现可以值得并行加速代码. ...
C#多线程学习(四) 多线程的自动管理(线程池) (转载系列)——继续搜索引擎研究...
在多线程的程序中,经常会出现两种情况: 一种情况: 应用程序中,线程把大部分的时间花费在等待状态,等待某个事件发生,然后才能给予响应这一般使用ThreadPo ...
python学习四(处理数据)
python学习四(处理数据) head first python中的一个数据处理的例子有四个U10选手的600米成绩,请取出每个选手跑的最快的3个时间.以下是四位选手的9次成绩 James 2-3 ...
PyTorch框架学习四——计算图与动态图机制
PyTorch框架学习四--计算图与动态图机制一.计算图二.动态图与静态图三.torch.autograd 1.torch.autograd.backward() 2.torch.autogra ...
Docker学习四：Docker 网络
前言本次学习来自于datawhale组队学习: 教程地址为: https://github.com/datawhalechina/team-learning-program/tree/master/ ...
（转）SpringMVC学习(四)——Spring、MyBatis和SpringMVC的整合
http://blog.csdn.net/yerenyuan_pku/article/details/72231763 之前我整合了Spring和MyBatis这两个框架,不会的可以看我的文章MyBa ...
算法学习四：算法性能分析理论基础——函数增长与渐进分析
算法学习四:算法性能分析理论基础--函数增长与渐进分析在算法性能分析过程中,特别是在算法运行效率分析中,我们经常使用渐渐分析法,它使我们在分析算法性能时不必纠结于不同硬件平台的差异性,着重考虑算法的 ...
Tensorflow学习四---高阶操作
Tensorflow学习四-高阶操作 Merge and split 1.tf.concat 拼接 a = tf.ones([4,32,8]) b = tf.ones([2,32,8]) print( ...
[jQuery学习系列四 ]4-Jquery学习四-事件操作
[jQuery学习系列四 ]4-Jquery学习四-事件操作前言: 今天看知乎偶然看到中国有哪些类似于TED的节目, 回答中的一些推荐我给记录下来了, 顺便也在这里贴一下: 一席云集听道推酷 ...
AVI音视频封装格式学习(四)——linux系统C语言AVI格式音视频封装应用
拖了很久的AVI音视频封装实例,花了一天时间终于调完了,兼容性不是太好,但作为参考学习使用应该没有问题.RIFF和AVI以及WAV格式,可以参考前面的一些文章.这里详细介绍将一个H264视频流和一个2 ...

opencl学习（四）

opencl学习（四）相关推荐

最新文章

热门文章