基于OpenCL的mean filter性能
1.对于一个标准的3*3 均值滤波,kernel代码如下:
使用buffer/image缓冲对象
__kernel void filter(__global uchar4* inputImage, __global uchar4* outputImage, uint N) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);int k = (N-1)/2;int n = N*N; //n*nif(x < k || y < k || x > width - k - 1 || y > height - k - 1) { outputImage[x + y * width] = inputImage[x + y * width];return; }uint4 finalcolor = (uint4)(0);int i,j;for(j = y - k; j <= y + k; j++) {for(i = x - k; i <= x + k; i++) { finalcolor = finalcolor + convert_uint4(inputImage[i + j * width]); } } outputImage[x + y * width] = convert_uchar4(finalcolor/n);}
__kernel void filterImg( image2d_t inputImage, __write_only image2d_t outputImage, uint N) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);uint4 temp = read_imageui(inputImage, imageSampler, (int2)(x,y));int k = (N-1)/2;int n = N*N; //n*nif(x < k || y < k || x > width - k - 1 || y > height - k - 1) { write_imageui(outputImage, (int2)(x,y), temp);return; }/* k*k area */ uint4 finalcolor = (uint4)(0);int i,j;for(j = y - k; j <= y + k; j++) {for(i = x - k; i <= x + k; i++) { finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(i,j)); } }finalcolor = finalcolor/n;write_imageui(outputImage, (int2)(x,y), finalcolor);}
对一个2048*2048的图像执行filter操作,
global work size = {2048, 2048, 1}, group work size = {16, 16}, 一般group work size应该为64的倍数,因为对于AMD显卡,wave是基本的硬件线程调度单位。
使用了6个GPRs,没有使用ScratchRegs,ScratchRregs是指用vedio meory来模拟GPR,但是线程执行的速度会大大降低,应尽量减少ScratchRegs的数量。
可以看到,使用image对象kernel执行时间要短,但奇怪的是各项性能参数都是buffer对象领先,除了alu busy和alu指令数目。
改为下面的kernel代码,性能会有所提高
__kernel void filter(__global uchar4* inputImage, __global uchar4* outputImage, uint N) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);if(x < 1 || y < 1 || x > width - 2 || y > height - 2) { outputImage[x + y * width] = inputImage[x + y * width];return; }uint4 finalcolor = (uint4)(0);finalcolor = finalcolor + convert_uint4(inputImage[x-1+( y-1) * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+( y-1) * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+1+( y-1) * width]); finalcolor = finalcolor + convert_uint4(inputImage[x-1+y * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+y * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+1+y * width]); finalcolor = finalcolor + convert_uint4(inputImage[x-1+( y+1) * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+( y+1) * width]); finalcolor = finalcolor + convert_uint4(inputImage[x+1+( y+1) * width]);outputImage[x + y * width] = convert_uchar4(finalcolor/9);} __kernel void filter1(__global uchar4* inputImage, __global uchar4* outputImage, uint N) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);int k = (N-1)/2;int n = N*N; //n*nif(x < k || y < k || x > width - k - 1 || y > height - k - 1) { outputImage[x + y * width inputImage[x + y * width];return; }// if(x==209 && y ==243)//{// printf("final color:%d,%d,%d,%d\n", finalcolor.x, finalcolor.y, finalcolor.z,finalcolor.w);// }uint4 finalcolor = (uint4)(0);int i,j;for(j = y - k; j <= y + k; j++) {for(i = x - k; i <= x + k; i++) { finalcolor = finalcolor + convert_uint4(inputImage[i + j * width]); } } outputImage[x + y * width] = convert_uchar4(finalcolor/n);} __kernel void filterImg( image2d_t inputImage, __write_only image2d_t outputImage, uint N) {int x = get_global_id(0);int y = get_global_id(1);int width = get_global_size(0);int height = get_global_size(1);uint4 temp = read_imageui(inputImage, imageSampler, (int2)(x,y));if(x < 1 || y < 1 || x > width - 2 || y > height - 2) { write_imageui(outputImage, (int2)(x,y), temp);return; }/* k*k area */ uint4 finalcolor = (uint4)(0);finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x-1,y-1)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x,y-1)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x+1,y-1)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x-1,y)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x,y)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x+1,y)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x-1,y+1)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x,y+1)); finalcolor = finalcolor + read_imageui(inputImage, imageSampler, (int2)(x+1,y+1));finalcolor = finalcolor/9;write_imageui(outputImage, (int2)(x,y), finalcolor);}
基于OpenCL的mean filter性能相关推荐
- 基于OpenCL的深度学习工具:AMD MLP及其使用详解
from:http://www.csdn.net/article/2015-08-05/2825390 [编者按]深度学习是近年来迅速发展和突破的机器学习领域,具有非常广泛的应用前景.将服务器GPU应 ...
- 基于OpenCL的数字地形分析之坡度坡向提取
转自:http://blog.csdn.net/zhouxuguang236/article/details/23568875 又有一段时间没有发表博客了,可能最近工作有点忙.今天就把最近的学习和研究 ...
- linux提升nvme性能,基于SPDK的NVMe SSD性能评估指南
原标题:基于SPDK的NVMe SSD性能评估指南 一 通过fio工具测试磁盘性能 SPDK采用异步I/O(Asynchronous I/O)加轮询(Polling)的工作模式,通常与Kernel的异 ...
- 基于容积卡尔曼滤波(CubatureKalmam Filter, CKF)的车辆状态观测器 Carsim与Simulink联合 可生成C代码
基于容积卡尔曼滤波(CubatureKalmam Filter, CKF)的车辆状态观测器 Carsim与Simulink联合 可生成C代码 ?CKF算法使用子函数形式编程,在定义好状态方程和观测方程 ...
- web前端高级JavaScript - 彻底掌握基于HTTP网络层的“前端性能优化”
彻底掌握基于HTTP网络层的 "前端性能优化" 产品性能优化方案 HTTP网络层优化 代码编译层优化 webpack 代码运行层优化 html/css javascript vue ...
- Tair LDB基于Prefixkey的范围查找性能优化项目中期总结
"Tair LDB基于Prefixkey的范围查找性能优化"这个项目刚好进行了一个月,这一个月主要是熟悉项目.掌握项目和提出设计方案的过程,下面从几个方面总结下个人在该项目上所做的 ...
- matlab在风能领域,基于Matlab/Simulink的风力机性能仿真研究
能源研究与信息 第 22 卷 第 2 期 Energy Research and Information Vol. 22 No. 2 2006 收稿日期:2005-10-14 作者简介:高 平(198 ...
- 基于OpenCL的图像积分图算法改进
复杂的算法却未必低效,简单的算法往往要付出代价,这个代价可能很大.在opencl环境下编程,与我们在CPU上的传统编程思想有一些差异,这些差异看似微不足道,但往往是细节决定成功,就是这些看似微不足道的 ...
- 计算机系统的性能建模与设计 排队理论,基于排队论的UM-BUS总线性能建模与评价-计算机科学.PDF...
第 44 卷第6A 期 Vo l. 44 No. 6A 计算机科学 2017 年 6 月 CO岛1PUTER SCIENCE June 2017 基于排队论的 UM-BUS 总线性能建模与评估 张少楠 ...
最新文章
- 20120907学习笔记
- 网管分析数据包? 组合一下吧。
- activiti 如何获取下一步节点
- 信息学奥赛一本通 1079:计算分数加减表达式的值 | OpenJudge NOI 1.5 33
- Windows环境下多个tomcat启动方法
- 使用mvc时,在视图view中使用强类型视图,在web.config文件中添加命名空间namespace的引用不起作用,解决方法...
- Centos下安装minikube
- HIVE 的笛卡尔积on无条件
- [转载] python 运算符重载有什么用_Python运算符重载用法实例分析
- 【细胞分割】基于matlab GUI阙值+边缘+形态学+种子点图像分割【含Matlab源码 615期】
- hdu 1560 DNA sequence(IDA*)
- johnson 算法 贪心
- 怎么通过创新再造58
- 联想微型计算机a20,联想乐player A20
- DH算法 | Diffie-Hellman 密钥交换
- 思岚雷达rplidar S1配置调试全纪录
- 中国企业服务总线(ESB)市场趋势报告、技术动态创新及市场预测
- 戴德金--连续性和无理数--我自己做的中文翻译第1页
- 折纸测珠峰python程序_哪个大神能帮忙用python做一下这个程序,真的万分感谢
- 暗影精灵双系统(win10和Ubuntu16.04)安装+Cuda和tensorflow安装-深度学习环境配置
热门文章
- 【转】DICOM的常用Tag分类和说明!!!!
- 【转】自旋锁-SpinLock(.NET 4.0+)
- 【转】事务和锁机制是什么关系? 开启事务就自动加锁了吗?
- 一步步编写操作系统 22 硬盘操作方法
- java 引用被回收_java GC 静态List 如果没有引用会被回收吗
- RSA公私钥加解密方式-工具类
- *【CodeForces - 859C 】Pie Rules (博弈dp,时光倒流)
- 【HDU - 1045】Fire Net (dfs 或二分图)
- 12.Cross-Validation
- mysql 磁盘组_有效管理 ASM 磁盘组空间