opencl 学习kernel,work_item和workgroup，工作组，工作项的理解
https://zhuanlan.zhihu.com/p/558573914
https://www.cnblogs.com/biglucky/p/3755189.html
转载自：http://www.cmnsoft.com/wordpress/?p=1429
https://blog.csdn.net/liyuan123zhouhui/article/details/52850282
前几节我们一起学习了几个用ＯＰＥＮＣＬ完成任务的简单例子，从这节起我们将更详细的对ＯＰＥＮＣＬ进行一些“理论”学习。
kernel：是指一个用opencl c语言编写的、代表一个单一执行实例的代码单元。opencl c语言看起来跟Ｃ语言函数非常相像，都有一个参数列表“局部”变量定义和标准控制流结构。opencl术语中把这种kernel实例称为work-item(工作项)。但opencl kernel与c语方函数的区别在于其并行语义。
work_item：是定义在一个很大的并行执行空间中的一小部分。是并行操作中每一部分的实例化。通俗来说，可以理解为kernel里定义的执行函数。当kernel启动后会创建大量work_item来同时执行，以完成并行任务。work_item根所其数据结构大小可分为一维、二维和三维数据。work_item之是的运行是相互独立的，不同步的。
work_group：opencl将全局执行空间划分为大量大小相等的，一维、二维、三维的work_item集合，这个集合就是work_group。在work_group内部，各个work_item之间允许一定程度的通信。而有work_group保证并发执行来允许其内部的work_item之间的本地同步。
在实际编写内核中，要了解线程调度的维度数，work_group的大小是很重要的，这有利于我们优化编写的内核程序。opencl提供了一此非常有用的函数供我们调用(在内核中调用)。
uint get_work_dim() : 返回线程调度的维度数。
uint get_global_size(uint dimension) : 返回在所请求维度上work_item的总数。
uint get_global_id(uint dimension) : 返回在所请求的维度上当前work_item在全局空间中的索引。
uint get_local_size(uint dimension) : 返回在所请求的维度上work-group的大小。
uint get_local_id(uint dimension) : 返回在所请求的维度上，当前work_item在work_group中的索引。
uint get_number_groups(uint dimension) : 返回在所请求维度上work-group的数目，这个值等于get_global_size 除以 get_local_size。
uint get_group_id(uint dimension) : 返回在所请求的维度上当前wrok_group在全局空间中的索引。
关于使用这些函数，我们举一个之前学过的例子。在第７课《旋转变换（１）》中的内核程序中原文是这样的。
这里传递的width和height大小是一样的，表示图像数据长宽的大小。其实也就是维度上work_item的总数，我们可以把代码改成。
把所有的width和heigh全部改成get_global_size(0),程序运行结果是一样的。而且我们还可以少传递两个参数。节省空间，提高效率。大家看下以前的例子，看看那些代码我们还可以优化呢。
首先给出工作项，工作组的概念：
工作项：最简单的理解，一个循环中最里面的一次运算，称为一个工作项。
工作组：是由访问相同处理资源的工作项组成，其主要优势有：
a 工作组中的工作项可以访问高速内存（也叫局部内存）的同一块内存
b 工作组中的工作项可以通过栅栏（fence）和障碍（barrier）的方式来进行同步
处理单元：能够支持工作组的处理资源被称为处理单元。各个工作组都可以在单个处理单元上的执行，而各个处理单元一次只能够执行一个工作组。
有时，会需要在设备上对数据进行划分，形成工作组、工作项，个人的理解，之所以opencl中需要划分工作组，工作项，是因为：
a 工作项之间需要进行同步化处理，如果工作项之间需要进行同步化处理，那么，就需要对这些工作项进行工作组划分。因为只有工作项属于同一个工作组，才能够实现工作项同步化。所谓工作项的同步化，是由于，多个工作项之间可能会访问局部内存中的同一段数据，他们之间的执行顺序就可能得到不同的结果，甚至是错误。为了防止这种错误的产生，就需要在不同的工作项之间进行同步，工作项之间的同步有两种方法：第一种是栅栏(fence)和障碍（barrier），第二种是原子操作。
那么一个内核可以生成多少个相对应的工作项？工作项和工作组的数量是没有限制的。但如果一个设备只含有M个计算单元，每个工作组有N个工作项，那么任何时候，最多只能有MN个工作项来执行内核（如上面已经提到，一个计算单元一次只能够执行一个工作组）。
在单个设备上能够进行数据划分的函数只有一个：clEnqueueNDRangeKernel，这个函数的声明以及重要参数说明如下：
1 clEnqueueNDRangeKernel(cl_command_queue queue,cl_kernel kernel,cl_uint work_dims,const size_t *global_work_offset,const size_t *global_work_size,const size_t *local_work_size,cl_uint num_events,const cl_event *wait_list,cl_event *event) 2 3 work_dims：数据的维度数，其限定范围一般是1-3 4 global_work_offset：各个维度上 全局ID偏移量，说白了，就是各个维度上数据的起始点，比如二维数据，你希望的起始点是(1,3),就可以通过这个参数设置，一般我们的起始点是从(0,0)开始的 5 global_work_size：各个维度上的工作项的数量，注意，这里说的各个维度上的工作项的数量，指的是总体的，因为在下面还有一个参数会指定在一个工作组上各个维度上的工作项的数量，这个参数是： 6 local_work_size：各个维度上一个工作组中工作项的数量
创建工作组并不是必须的，如果参数local_work_size的取值被设置成NULL，opencl将分析决定如何在设备上的处理单元间分配工作项。
举例：
对于一个12x12的矩阵，将其划分为9个工作组，每个工作组是一个4x4的矩阵，有16个工作项，相应的设置：
1 cl_uint work_dims =2; //2维数据 2 size_t global_work_offset[2] = {0,0}; //从0，0开始 3 size_t global_work_size[2] = {12,12}; 4 size_t local_work_size[2] = {4,4}; 5 clEnqueueNDRangeKernel(command_q,kernel,work_dim,global_work_offset,global_work_size,local_work_size,0,NULL,NULL);
而对于另一个运算，比如，将一个8x8的图像与一个3x3的核进行卷积，放到一个6x6的图像中，如果不分组，可以直接将二维图像进行一维访问也可以：
1 size_t global_work_size[2] = {36};/6x6 2 size_t local_work_size[2] = {1}; 3 clEnqueueNDRangeKernel(command_q,kernel,1,NULL,global_work_size,local_work_size,0,NULL,NULL);
————————————————
版权声明：本文为CSDN博主「跬步达千里」的原创文章，遵循CC 4.0 BY-SA版权协议，转载请附上原文出处链接及本声明。
原文链接：https://blog.csdn.net/liyuan123zhouhui/article/details/52850282
