OpenCL 第10课:kernel,work_item和workgroup

转载自：http://www.cmnsoft.com/wordpress/?p=1429

前几节我们一起学习了几个用ＯＰＥＮＣＬ完成任务的简单例子，从这节起我们将更详细的对ＯＰＥＮＣＬ进行一些“理论”学习。

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在全局空间中的索引。

关于使用这些函数，我们举一个之前学过的例子。在第７课《旋转变换（１）》中的内核程序中原文是这样的。

__kernel void rotation(__global int* A,
                    __global int* B,
                    int width,
                    int height,
                    float sinangle,
                    float cosangle)
{
    //获取索引号，这里是二维的，所以可以取两个
    //否则另一个永远是0
    int col = get_global_id(0);
    int row = get_global_id(1);
 
    //计算图形中心点
    float cx = ((float)width)/2;
    float cy = ((float)height)/2;
 
    int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
    int ny = (int)(cy + (-1*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy));
 
    //边界检测
    if(nx>=0 && nx<width && ny>=0 && ny<height)
    {
        B[nx + ny*width] = A[col + row*width];
    }
 
}

这里传递的width和height大小是一样的，表示图像数据长宽的大小。其实也就是维度上work_item的总数，我们可以把代码改成。

__kernel void rotation(__global int* A,
                    __global int* B,
                    int width,
                    int height,
                    float sinangle,
                    float cosangle)
{
    //获取索引号，这里是二维的，所以可以取两个
    //否则另一个永远是0
    int col = get_global_id(0);
    int row = get_global_id(1);
 
    //计算图形中心点
    float cx = ((float)get_global_size(0))/2;
    float cy = ((float)get_global_size(0))/2;
 
    int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
    int ny = (int)(cy + (-1*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy));
 
    //边界检测
    if(nx>=0 && nx<get_global_size(0) && ny>=0 && ny<get_global_size(0))
    {
        B[nx + ny*get_global_size(0)] = A[col + row*get_global_size(0)];
    }
 
}

把所有的width和heigh全部改成get_global_size(0),程序运行结果是一样的。而且我们还可以少传递两个参数。节省空间，提高效率。大家看下以前的例子，看看那些代码我们还可以优化呢。