OpenCL概述 续篇
来源:互联网 发布:路由器劫持软件 编辑:程序博客网 时间:2024/06/05 02:14
Kernel对象:
Kernel就是在程序代码中的一个函数,这个函数能在OpenCL设备上执行。一个Kernel对象就是kernel函数以及其相关的输入参数。
Kernel对象通过程序对象以及指定的函数名字创建。注意:函数必须是程序源代码中存在的函数。
运行时编译:
在运行时,编译程序和创建kernel对象是有时间开销的,但这样比较灵活,能够适应不同的OpenCL硬件平台。程序动态编译一般只需一次,而Kernel对象在创建后,可以反复调用。
创建Kernel后,运行Kernel之前,我们还要为Kernel对象设置参数。我们可以在Kernel运行后,重新设置参数再次运行。
arg_index指定该参数为Kernel函数中的第几个参数(比如第一个参数为0,第二个为1,…)。内存对象和单个的值都可以作为Kernel参数。下面是2个设置Kernel参数的例子:
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_iImage);
clSetKernelArg(kernel, 1, sizeof(int), (void*)&a);
在Kernel运行之前,我们先看看OpenCL中的线程结构:
大规模并行程序中,通常每个线程处理一个问题的一部分,比如向量加法,我们会把两个向量中对应的元素加起来,这样,每个线程可以处理一个加法。
下面我看一个16个元素的向量加法:两个输入缓冲A、B,一个输出缓冲C
在这种情况下,我们可以创建一维的线程结构去匹配这个问题。
每个线程把自己的线程id作为索引,把相应元素加起来。
OpenCL中的线程结构是可缩放的,Kernel的每个运行实例称作WorkItem(也就是线程),WorkItem组织在一起称作WorkGroup,OpenCL中,每个Workgroup之间都是相互独立的。
通过一个global id(在索引空间,它是唯一的)或者一个workgroup id和一个work group内的local id,我就能标定一个workitem。
在kernel函数中,我们能够通过API调用得到global id以及其他信息:
get_global_id(dim)
get_global_size(dim)
这两个函数能得到每个维度上的global id。
get_group_id(dim)
get_num_groups(dim)
get_local_id(dim)
get_local_size(dim)
这几个函数用来计算group id以及在group内的local id。
get_global_id(0) = column, get_global_id(1) = row
get_num_groups(0) * get_local_size(0) == get_global_size(0)
OpenCL内存模型
OpenCL的内存模型定义了各种各样内存类型,各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的,比如从host memory到device memory,从global memory到local memory等等。
WorkGroup被映射到硬件的CU上执行(在AMD 5xxx系列显卡上,CU就是simd,一个simd中有16个pe,或者说是stream core),OpenCL并不提供各个workgroup之间的一致性,如果我们需要在各个workgroup之间共享数据或者通信之类的,要自己通过软件实现。
Kernel函数的写法
每个线程(workitem)都有一个kenerl函数的实例。下面我们看下kernel的写法:
每个Kernel函数都必须以__kernel开始,而且必须返回void。每个输入参数都必须声明使用的内存类型。通过一些API,比如get_global_id之类的得到线程id。
内存对象地址空间标识符有以下几种:
__global – memory allocated from global address space
__constant – a special type of read-only memory
__local – memory shared by a work-group
__private – private per work-item memory
__read_only/__write_only – used for images
Kernel函数参数如果是内存对象,那么一定是__global,__local或者constant。
运行Kernel
首先要设置线程索引空间的维数以及workgroup大小等。
我们通过函数clEnqueueNDRangeKerne把Kernel放在一个队列里,但不保证它马上执行,OpenCL driver会管理队列,调度Kernel的执行。注意:每个线程执行的代码都是相同的,但是它们执行数据却是不同的。
该函数把要执行的Kernel函数放在指定的命令队列中,globald大小(线程索引空间)必须指定,local大小(work group)可以指定,也可以为空。如果为空,则系统会自动根据硬件选择合适的大小。event_wait_list用来选定一些events,只有这些events执行完后,该kernel才可能被执行,也就是通过事件机制来实现不同kernel函数之间的同步。
当Kernel函数执行完毕后,我们要把数据从device memory中拷贝到host memory中去。
释放资源:
大多数的OpenCL资源都是指针,不使用的时候需要释放掉。当然,程序关闭的时候这些对象也会被自动释放掉。
释放资源的函数是:clRelase{Resource} ,比如: clReleaseProgram(), clReleaseMemObject()等。
错误捕捉:
如果OpenCL函数执行失败,会返回一个错误码,一般是个负值,返回0则表示执行成功。我们可以根据该错误码知道什么地方出错了,需要修改。错误码在cl.h中定义,下面是几个错误码的例子.
CL_DEVICE_NOT_FOUND -1
CL_DEVICE_NOT_AVAILABLE -2
CL_COMPILER_NOT_AVAILABLE -3
CL_MEM_OBJECT_ALLOCATION_FAILURE -4
…
下面是一个OpenCL机制的示意图
程序模型
数据并行:work item和内存对象元素之间是一一映射关系;workgroup可以显示指定,也可以隐式指定。
任务并行:kernel的执行独立于线程索引空间;用其他方法表示并行,比如把不同的任务放入队列,用设备指定的特殊的向量类型等等。
同步:workgroup内work item之间的同步;命令队列中不同命令之间的同步。
完整代码如下:
也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下载完整版本。
原文作者:迈克老狼
- OpenCL概述 续篇
- OpenCL概述 续篇(Introduction to OpenCL)
- OpenCL概述
- OpenCL概述 (Introduction to OpenCL)
- OpenCL
- OpenCL
- OpenCL
- OpenCL
- opencl
- OpenCL
- OpenCL
- OpenCL
- OpenCl
- OpenCL
- 回国续篇
- DynamicQueryable续篇
- OpenCL: OpenCL架构
- OpenCL: OpenCL的shader
- 【裸奔吧linux】ubuntu 个人收藏软件列表
- 漫谈协同过滤推荐算法(一)
- 10_JSP 隐式对象
- Java微信公众号开发:网页授权获取用户基本信息
- c++数字和字符串的转换
- OpenCL概述 续篇
- 11_JSP 客户端请求
- 06:月度开销OJ
- SlidingPaneLayout双栏滑动
- 树——populating-next-right-pointers-in-each-node(层序遍历变型)
- 12_JSP 服务器响应
- Java内部类的使用小结
- GPU架构
- Android/linux 平台下使用 i2c-tools调试IIC方法