转载自:http://www.cmnsoft.com/wordpress/?p=1429
前几节我们一起学习了几个用OPENCL完成任务的简单例子,从这节起我们将更详细的对OPENCL进行一些“理论”学习。
kernel:是指一个用opencl c语言编写的、代表一个单一执行实例的代码单元。opencl c语言看起来跟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在全局空间中的索引。
关于使用这些函数,我们举一个之前学过的例子。在第7课《旋转变换(1)》中的内核程序中原文是这样的。
__kernel void rotation(__global int* A,
__global int* B,
int width,
int height,
float sinangle,
float cosangle)
{
//获取索引号,这里是二维的,所以可以取两个
//否则另一个永远是0
int col = get_global_id();
int row = get_global_id(); //计算图形中心点
float cx = ((float)width)/;
float cy = ((float)height)/; int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
int ny = (int)(cy + (-*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy)); //边界检测
if(nx>= && nx<width && ny>= && 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();
int row = get_global_id(); //计算图形中心点
float cx = ((float)get_global_size())/;
float cy = ((float)get_global_size())/; int nx = (int)(cx + cosangle * ((float)col-cx) + sinangle * ((float)row-cy));
int ny = (int)(cy + (-*sinangle) * ((float)col-cx) + cosangle * ((float)row-cy)); //边界检测
if(nx>= && nx<get_global_size() && ny>= && ny<get_global_size())
{
B[nx + ny*get_global_size()] = A[col + row*get_global_size()];
} }
把所有的width和heigh全部改成get_global_size(0),程序运行结果是一样的。而且我们还可以少传递两个参数。节省空间,提高效率。大家看下以前的例子,看看那些代码我们还可以优化呢。