opencl
文章平均质量分 70
10km
这个作者很懒,什么都没留下…
展开
-
OpenCL Installable Client Driver (ICD) Loader编译
ICD是什么?OpenCL Installable Client Driver (ICD) Loader是实现OpenCL应用程序与各硬件厂商提供的OpenCL驱动(platform)之间隔离的中间库。从OpenCL 1.2开始,OpenCL提供了一个ICD扩展(cl_khr_icd),它允许不同厂商的多个OpenCL驱动(platform)共存于一个主机系统,应用程序可以通过调用clIcdGetP原创 2016-01-11 10:38:09 · 5937 阅读 · 1 评论 -
opencl:c++接口(cl.hpp)利用cl::LocalSpaceArg设置__local 参数
当我们需要在kernel中使用local memory数组的时候,有两种方式定义local 数组 第一种,编译期静态定义,这是比较普通的使用方式,如下代码,这种方式,在编译期就分配了local 数组的大小。#define LOCAL_ARRAY_SIZE 64 // LOCAL_ARRAY_SIZE 可以通过编译选项-D在编译的时候定义__kernel void test_kernel( ){原创 2016-05-16 16:33:37 · 2376 阅读 · 1 评论 -
opencl/msvc:kernel因为指针对齐方式(alignment)造成向量类型(vector data type)读写异常
opencl knernel中对全局内存(__global)向量类型数据的读写有两种方式, 一种是直接用=操作符赋值,一种则是通过vstoren/vloadn函数来实现向量数据读写。 =操作符赋值方式使用简单,但在msvc下以CL_MEM_USE_HOST_PTR模式向kernel传递数据时如果数据对齐方式不对,会造成kernel运行时异常。本文就是针对这种情况分析原因并提出解决方案。原创 2016-04-19 12:03:55 · 3179 阅读 · 1 评论 -
opencl:慎用-cl-opt-disable选项编译kernel(可能会导致一些无法解释的问题)
在编译opencl kernel代码时,有一个编译选项-cl-opt-disable。根据opencl 官网的原文描述,使用这个选项可以关闭所有的代码优化,便于调试程序。(默认情况下,编译优化选项是打开的) 参见clBuildProgram 但是今天为了调试kernel代码,实际使用这个选项编译kernel却发现,使用这个选项就是坑。使用之后,kernel参数传递都不正常了。 下面这是个很简单原创 2016-04-21 15:00:01 · 2522 阅读 · 1 评论 -
opencl:原子命令实现自旋锁(spinlock)的使用限制
原子命令很重要的用途就是互斥(mutexes)。互斥保证了每次只有一个work-item能访问数据。opencl也支持原子命令,在opencl最初始的版本1.0,原子命令是作为扩展功能(opencl extensions)来提供的(参见cl_khr_global_int32_base_atomics,cl_khr_global_int32_extended_atomics)。到opencl1.2以后原创 2016-05-29 16:11:10 · 2549 阅读 · 1 评论 -
opencl:提高kernel有效并发约束(Effective concurrency constraint)的一点心得
下面是一个简单的kernel函数,从integ_count_mat矩阵中用vload8函数取出A1,A2,A3,A4四个向量执行A4+A1-A2-A3,结果存入density_mat,代码中只用到了一个向量类型的变量sum。__kernel void object_density_filter( matrix_info_cl im_info , const __原创 2016-05-09 14:51:17 · 2965 阅读 · 1 评论 -
opencl:获取每个计算单元(CU)中处理元件(PE)的数目
OpenCL 平台模型的定义如下图。模型中有一个主机,并且有一个或多个OpenCL 设备与其相连。每个OpenCL 设备可划分成一个或多个计算单元(CU),每个计算单元又可划分 成一个或多个处理元件(PE)。设备上的计算是在处理元件中进行的。 OpenCL 应用程序会按照主机平台的原生模型在这个主机上运行。主机上的OpenCL 应用程 序提交命令(command queue)给设备中的处理元原创 2016-05-23 11:01:25 · 7854 阅读 · 2 评论 -
基于OpenCL的图像积分图算法改进
简单往要付出代价,这个代价可能很大,在opencl环境下编程,与我们在CPU上的传统编程思想有一些差异,这些差异看似微不足道,但往往细节决定成功,就是这些看似微不足道的差异导致同一种算法在GPU和CPU运行效果有着巨大的差别 之前写过一篇文章《基于OpenCL的图像积分图算法实现》介绍了opencl中积分图算法的基本原理(不了解积分图概念的朋友可以先参考这篇文章),并基于这个基本原理提供了kern原创 2016-06-08 10:25:01 · 5210 阅读 · 14 评论 -
基于OpenCL的图像积分图算法实现
积分图的概念图像积分图算法在图像特征检测中有着比较广泛的应用,主要用于规则区域特征值的计算。 积分图的概念可用下图表示: 坐标A(x,y)的积分图是其左上角的所有像素之和(图中的阴影部分)。定义为: 在上图中,A(x,y)表示点(x,y)的积分图;s(x,y)表示点(x,y)的y方向的所有原始图像之和。积分图算法在CPU上的串行实现在CPU上串行实现积分图计算的典型代码如下: /*原创 2016-03-12 16:51:50 · 4004 阅读 · 2 评论 -
opencl::kernel中获取local memory size
在OpenCL设备中一个workgroup中的所有work-item可以共用本地内存(local memory),在OpenCL kernal编程中,合理的利用local memory,可以提升系统的整体效率。 但是,根据OpenCL的标准,不论在kernel代码的编译期还是运行时,kernel程序在不借助主机端程序的帮助下,是无法知道当前设备(device)的local memory容量的。也就原创 2016-03-04 15:22:39 · 3875 阅读 · 0 评论 -
opencl:异步复制函数的注意事项(async_work_group_copy/async_work_group_strided_copy)
OpenCL中的内置函数async_work_group_copy和async_work_group_strided_copy用于实现全局内存(global memory)和本地内存(local memory)之间的异步数据复制,在某些情况下,使用异步复制(async copy)的方式在全局内存和本地内存之间复制数据比直接赋值的方式要方便。 下面是async_work_group_copy的函数说原创 2016-06-11 13:21:05 · 3386 阅读 · 2 评论 -
opencl:C++实现双线性插值图像缩放
用OpenCL实现图像缩放代码是比较简单的,因为OpenCL本身就支持双线性插值 下面是kernel代码(从Mali OpenCL SDK 抄来的:/samples/image_scaling/assets/image_scaling.cl) 非常简单只有4行// 定义采样器 // CLK_NORMALIZED_COORDS_TRUE指定使用归一化坐标// CLK_ADDRESS_CLAMP原创 2016-02-27 16:20:24 · 8457 阅读 · 4 评论 -
opencl:clEnqueueNDRangeKernel执行报错CL_OUT_OF_RESOURCES的一种情况
我的电脑上之前的显卡比较老并不支持opencl,所以我之前开发时opencl代码其实都是在CPU上跑的,现在所有的代码都调试通过了,决定装块新显卡用于程序的性能测试。 今天显卡到了,装上之后运行程序,clEnqueueNDRangeKernel在执行下面的kernel时报错:CL_OUT_OF_RESOURCES。__kernel void prefix_sum_col_and_transpose原创 2016-05-03 16:02:00 · 5555 阅读 · 1 评论 -
opencl:一个关于向量赋值的异常
在项目中,有一个下面这样的数据结构,storage保存是个float4类型的数组。typedef struct _detected_objects_buffer { cl_float4 storage[MAX_DETECTED_OBJECT_NUM]; cl_int detected_num; kernel_error status;}detected_objects_b原创 2016-04-17 11:28:17 · 1994 阅读 · 1 评论 -
C++代码设计:向Java借鉴Builder模式塈OpenCL内核代码编译
Builder模式所谓的builder模式是指在设计Java代码时,当方法调用的参数过多的时候,可以用builder模式将所有参数封装在一个类中,然后将这个类的实例做为参数传递给方法。这样以来方法只需要接收一个类参数,就能获取所有想要的参数,尤其是对于多个类似方法,都需要差不多相同的参数的情况下,这种设计就更加有效率,可以减少方法调用的复杂度,减少出错的机会,如果你还不懂什么叫builder模式,这原创 2016-03-03 12:30:22 · 1423 阅读 · 1 评论 -
opencl:改造C++接口增加对内存编译(compile)的支持
OpenCL 1.2以前的标准(1.0,1.1),只支持单个源文件编译成可执行程序(Executable Program),所以只提供了clBuildProgram函数。 从OpenCL 1.2以后,可以将complie/link两个动作分开,增加了clCompileProgram, clLinkProgram函数,允许将多个源码编译成一个可执行程序,clCompileProgram将一段内核代码原创 2016-03-03 16:38:39 · 2099 阅读 · 3 评论 -
OpenCL: kernel中的向量关系运算符和等价运算符(>,<,>=,<=,==,!=)
opencl的kernel编程语言是C99标准的一个子集,在C99的基础上opencl增加了向量数据类型(Scalar Data Types):charn,ucharn,shortn,ushortn,intn,uintn,longn,ulongn,floatn,doublen(n=2,4,8,16)。 这些向量类型与基础的标量类型(Vector Data Types)一样支持各种算术和逻辑运算符。原创 2016-03-20 10:51:39 · 2491 阅读 · 0 评论 -
opencl:C++ 利用cl::make_kernel简化kernel执行代码
上一篇博客《opencl:C++实现双线性插值图像缩放》中介绍了简单的图像缩放函数 代码片段如下,可以看到,为了执行kernel,先要设置kernel参数,然后调用enqueueNDRangeKernel执行kernel。/* 缩放图像(双线性插值) 返回缩放后的图像数据*/gray_matrix_cl gray_matrix_cl::zoom(size_t dst_width, size_t原创 2016-02-29 18:19:44 · 4195 阅读 · 2 评论 -
opencl:cl::make_kernel的进化
我之前的一篇博客《opencl:C++ 利用cl::make_kernel简化kernel执行代码》详细说明了如何使用OpenCL C++接口(cl.hpp)提供cl::make_kernel算子来简化kernel执行代码。/* 缩放图像(双线性插值) */gray_matrix_cl gray_matrix_cl::zoom(size_t dst_width, size_t dst_height原创 2016-03-10 12:18:19 · 2569 阅读 · 1 评论 -
opencl: C++ 接口(cl.hpp)创建kernel
OpenCL不仅提供了标准C接口,同时提供C++的接口(cl.hpp),其实就是基于C接口的进一步封装。有了这个C++接口,对于C++项目来说,就大大提高了使用的便利性,本人涉及的这个项目对OpenCL的调用全部都是基于OpenCL的C++接口来完成的。本文讲述如何用OpenCL 1.2的C++接口来从cl原文件创建kernel。在cl.hpp中对cl_kernel被封装成了cl::Kernel对象原创 2016-02-27 15:21:37 · 6445 阅读 · 20 评论 -
opencl:C++接口配置抛出异常
在使用OpenCL的C++接口(cl.hpp)时,默认情况下所有的调用出错都是以错误码的形式返回的,如果希望opencl调用出错时以exception形式抛出,则可以像下面这样在#include语句之前加入__CL_ENABLE_EXCEPTIONS定义// 设置当opencl出错时抛出异常#ifndef __CL_ENABLE_EXCEPTIONS#define __CL_ENABLE_EXC原创 2016-02-27 16:38:20 · 1239 阅读 · 0 评论 -
opencl C++接口: 关于CL_KERNEL_FUNCTION_NAME的一个坑
我的项目中所有的kernel在程序初始化时就被编译生成了,存放在一个std::unordered_map类型的map表中(kernel name为key),以后程序需要调用的时候,就通过kernel name来获取指定的cl::Kernel对象。 建这个表的时候,要创建cl::Kernel。常用的创建cl::Kernel的途径有两个:cl::Pro原创 2016-03-13 12:32:54 · 2500 阅读 · 1 评论 -
opencl:C++11下使用别名(x,y,z,hi,lo...)访问vector类型(cl_int2,cl_long16...)的元素
在gcc(5.2.0)下使用C++11写opencl的程序时,发现无法像内核代码一样对cl_int2这样的向量(vector)类型用pos.x,pos.y这样的方式来访问向量元素,只能用pos.s[0]这种数组访问的方式。这是为什么?这是platform.h中cl_int2的定义,可以看出,虽然代码中有,x,y名字定义,但编译开关__CL_HAS_ANON_STRUCT__导致这部分代码是灰的原创 2016-04-10 10:57:12 · 5419 阅读 · 0 评论 -
C++11:模板实现opencl向量类型(cl_intn,cl_floatn...)的简单运算符(+,-)重载及length,distance函数
opencl内核支持的所有向量数据类型(intn,floatn,doublen….)在主机端都有对应的类型,区别是加了前缀cl_,比如int4对应的主机端类型是cl_int4。 我们知道,在opencl内核代码中,向量类型(vector data type)的数据可以像普通标量类型(scala data type)一样,用各种算术和逻辑运算符进行操作。 比如:int4 p1=int4(4,2,0原创 2016-04-11 15:10:48 · 4027 阅读 · 2 评论 -
opencl:kernel中两种向量类型转换(convert_T,as_typen)的主要区别
opencl kernel中向量类型转换分为两种方式,explicit conversions和reinterpreting type,中文可以分别直译为"显式转换"和"重新解释类型"。本文讨论这两种类型转换的区别。原创 2016-04-17 10:53:09 · 3820 阅读 · 1 评论 -
OpenCL ICD Loader运行测试暨解决报错:ERROR: App log and stub log differ.
上一篇博文《OpenCL Installable Client Driver (ICD) Loader编译》详细描述了如何编译OpenCL ICD Loader。OpenCL ICD Loader自带了测试程序,成功编译后可以根据源码根目录下README.txt的说明运行测试程序来验证Loader是否可以正常工作: == Running ICD Test == ICD Test can b原创 2016-01-11 15:54:21 · 2556 阅读 · 2 评论