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性能相关推荐

  1. 基于OpenCL的深度学习工具:AMD MLP及其使用详解

    from:http://www.csdn.net/article/2015-08-05/2825390 [编者按]深度学习是近年来迅速发展和突破的机器学习领域,具有非常广泛的应用前景.将服务器GPU应 ...

  2. 基于OpenCL的数字地形分析之坡度坡向提取

    转自:http://blog.csdn.net/zhouxuguang236/article/details/23568875 又有一段时间没有发表博客了,可能最近工作有点忙.今天就把最近的学习和研究 ...

  3. linux提升nvme性能,基于SPDK的NVMe SSD性能评估指南

    原标题:基于SPDK的NVMe SSD性能评估指南 一 通过fio工具测试磁盘性能 SPDK采用异步I/O(Asynchronous I/O)加轮询(Polling)的工作模式,通常与Kernel的异 ...

  4. 基于容积卡尔曼滤波(CubatureKalmam Filter, CKF)的车辆状态观测器 Carsim与Simulink联合 可生成C代码

    基于容积卡尔曼滤波(CubatureKalmam Filter, CKF)的车辆状态观测器 Carsim与Simulink联合 可生成C代码 ?CKF算法使用子函数形式编程,在定义好状态方程和观测方程 ...

  5. web前端高级JavaScript - 彻底掌握基于HTTP网络层的“前端性能优化”

    彻底掌握基于HTTP网络层的 "前端性能优化" 产品性能优化方案 HTTP网络层优化 代码编译层优化 webpack 代码运行层优化 html/css javascript vue ...

  6. Tair LDB基于Prefixkey的范围查找性能优化项目中期总结

    "Tair LDB基于Prefixkey的范围查找性能优化"这个项目刚好进行了一个月,这一个月主要是熟悉项目.掌握项目和提出设计方案的过程,下面从几个方面总结下个人在该项目上所做的 ...

  7. matlab在风能领域,基于Matlab/Simulink的风力机性能仿真研究

    能源研究与信息 第 22 卷 第 2 期 Energy Research and Information Vol. 22 No. 2 2006 收稿日期:2005-10-14 作者简介:高 平(198 ...

  8. 基于OpenCL的图像积分图算法改进

    复杂的算法却未必低效,简单的算法往往要付出代价,这个代价可能很大.在opencl环境下编程,与我们在CPU上的传统编程思想有一些差异,这些差异看似微不足道,但往往是细节决定成功,就是这些看似微不足道的 ...

  9. 计算机系统的性能建模与设计 排队理论,基于排队论的UM-BUS总线性能建模与评价-计算机科学.PDF...

    第 44 卷第6A 期 Vo l. 44 No. 6A 计算机科学 2017 年 6 月 CO岛1PUTER SCIENCE June 2017 基于排队论的 UM-BUS 总线性能建模与评估 张少楠 ...

最新文章

  1. 20120907学习笔记
  2. 网管分析数据包? 组合一下吧。
  3. activiti 如何获取下一步节点
  4. 信息学奥赛一本通 1079:计算分数加减表达式的值 | OpenJudge NOI 1.5 33
  5. Windows环境下多个tomcat启动方法
  6. 使用mvc时,在视图view中使用强类型视图,在web.config文件中添加命名空间namespace的引用不起作用,解决方法...
  7. Centos下安装minikube
  8. HIVE 的笛卡尔积on无条件
  9. [转载] python 运算符重载有什么用_Python运算符重载用法实例分析
  10. 【细胞分割】基于matlab GUI阙值+边缘+形态学+种子点图像分割【含Matlab源码 615期】
  11. hdu 1560 DNA sequence(IDA*)
  12. johnson 算法 贪心
  13. 怎么通过创新再造58
  14. 联想微型计算机a20,联想乐player A20
  15. DH算法 | Diffie-Hellman 密钥交换
  16. 思岚雷达rplidar S1配置调试全纪录
  17. 中国企业服务总线(ESB)市场趋势报告、技术动态创新及市场预测
  18. 戴德金--连续性和无理数--我自己做的中文翻译第1页
  19. 折纸测珠峰python程序_哪个大神能帮忙用python做一下这个程序,真的万分感谢
  20. 暗影精灵双系统(win10和Ubuntu16.04)安装+Cuda和tensorflow安装-深度学习环境配置

热门文章

  1. 【转】DICOM的常用Tag分类和说明!!!!
  2. 【转】自旋锁-SpinLock(.NET 4.0+)
  3. 【转】事务和锁机制是什么关系? 开启事务就自动加锁了吗?
  4. 一步步编写操作系统 22 硬盘操作方法
  5. java 引用被回收_java GC 静态List 如果没有引用会被回收吗
  6. RSA公私钥加解密方式-工具类
  7. *【CodeForces - 859C 】Pie Rules (博弈dp,时光倒流)
  8. 【HDU - 1045】Fire Net (dfs 或二分图)
  9. 12.Cross-Validation
  10. mysql 磁盘组_有效管理 ASM 磁盘组空间