基于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操作,

image

image

 

image

image

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);}

image

image

image

image

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/448820.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

Docker 实战:编写 Dockerfile

一、编译镜像 前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 1. 编译镜像 Dockerfile类似于Makfile&#xff0c;用户使用docker build就可以编译镜像&#xff0c;使用该命令可以设置编译…

dubbo-环境搭建,实现一个简单地dubbo实例(附github地址)

一、建立maven模块和provider、consumer、service子模块&#xff0c;其中service是开发接口的模块 建立一个maven模块&#xff0c;不选择样板&#xff0c;直接next知道完成&#xff0c;建立三个子模块,建立完后发现各个模块的java目录不是源目录 右键——>make Directory as…

static 二次理解

当api底层用到static修饰的话&#xff0c;因为是类的&#xff0c;此容器中只有一份转载于:https://blog.51cto.com/jiaxiaoxu/2394844

AMD 5XXX 系列显卡的 peak bandwidth计算

在ATI Stream Computing Programming Guide中&#xff0c;例举了AMD 5系列显卡的参数信息。 我比较关注其中Peak bandwidths的计算&#xff0c;以便在opencl程序测试bandwidth利用率。 下面&#xff0c;我以5870为例&#xff0c;探讨一下如何计算得到这些结果&#xff1a; L1 c…

Docker : Dockerfile 定制镜像

使用 Dockerfile 定制镜像 前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 镜像的定制实际上就是定制每一层所添加的配置、文件。如果我们可以把每一层修改、安装、构建、操作的命令都写…

动态规划 最长上升子序列

题意&#xff1a;给出一个序列&#xff0c;求它的最长上升子序列的长度 题目链接&#xff1a;https://ac.nowcoder.com/acm/problem/26156 输入:n代表长度&#xff0c;然后是一个字符串 分析&#xff1a;用dp[i]表示长度为i1的上升子序列末尾元素的最小值&#xff08;一开始初始…

解说redis中如何实现高可用

redis中为了实现高可用&#xff08;High Availability&#xff0c;简称HA&#xff09;&#xff0c;采用了如下两个方式&#xff1a;主从复制数据。采用哨兵监控数据节点的运行情况&#xff0c;一旦主节点出现问题由从节点顶上继续进行服务。主从复制redis中主从节点复制数据有全…

OpenCL memory object 之 Global memory (1)

这篇日志是学习AMD OpenCL文档时候的总结。 OpenCL用memory object在host和device之间传输数据&#xff0c;memory object由runtime&#xff08;运行库&#xff0c;driver的一部分&#xff09;来管理。 OpenCL中的内存对象包括buffer以及image&#xff0c;buffer是一维数据元素…

Docker: dockerfile 使用介绍

Docker简介 前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 Docker项目提供了构建在Linux内核功能之上&#xff0c;协同在一起的的高级工具。其目标是帮助开发和运维人员更容易地跨系统跨…

【Hello CSS】第六章-文档流与排版

作者&#xff1a;陈大鱼头github&#xff1a; KRISACHAN正常流 什么是“正常流”&#xff1f; 其实就是我们日常所说的“文档流”。 在W3C官方文档里对应的是“normal flow”。 正常流的盒子属于格式化上下文(FC)&#xff0c;在CSS2.2中可以是表格、块或内联。 在CSS3中引入了f…

创建型模式---工厂模式

工厂模式 在工厂设计模式中&#xff0c;客户端可以请求一个对象&#xff0c;而无需要知道这个对象来自哪里&#xff0c;也就是使用哪个类来生成这个对象。工厂背后的思想是简化对象的创建。与客户端自己基于类实例化直接创建对象相比&#xff0c;基于一个中心化函数来实现&…

OpenCL memory object 之 Global memory (2)

当我们用clCreateBuffer, clCreateImage创建OpenCL memory object时候&#xff0c;我们需要输入一个flag参数&#xff0c;这个参数决定memory object的位置。 cl_mem clCreateBuffer (cl_context context, cl_mem_flags flags, size_t size, void *host_ptr, cl_int *errc…

数据结构进阶篇-跳表

大家想必都知道&#xff0c;数组和链表的搜索操作的时间复杂度都是O(N)的&#xff0c;在数据量大的时候是非常耗时的。对于数组来说&#xff0c;我们可以先排序&#xff0c;然后使用二分搜索&#xff0c;就能够将时间复杂度降低到O(logN)&#xff0c;但是有序数组的插入是一个O…

查看本机ssh公钥,生成公钥

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 查看ssh公钥方法&#xff1a; 1.通过命令窗口&#xff1a;打开你的git bash 窗口&#xff0c;进入.ssh目录&#xff1a;cd ~/.ssh&…

如何实现动态水球图 --》 echars结合echarts-liquidfill实现

1&#xff09;项目中作为项目依赖&#xff0c;安装到项目当中(注意必须要结合echars) npm install echarts vue-echarts --save npm install echarts-liquidfill --save 2&#xff09;在需要使用水晶球的组件里引入liquidFill.js import echarts-liquidfill/src/liquidFill.js;…

OpenCL memory object 之选择传输path

对应用程序来说&#xff0c;选择合适的memory object传输path可以有效提高程序性能。 下面先看一写buffer bandwidth的例子&#xff1a; 1. clEnqueueWriteBuffer()以及clEnqueueReadBuffer() 如果应用程序已经通过malloc 或者mmap分配内存&#xff0c;CL_MEM_USE_HOST_PTR是个…

struts入门超详细

https://blog.csdn.net/yerenyuan_pku/article/details/52652262转载于:https://www.cnblogs.com/liuna369-4369/p/10870873.html

RabbitMQ 从入门到精通 (一)

目录 1. 初识RabbitMQ2. AMQP3.RabbitMQ的极速入门4. Exchange(交换机)详解4.1 Direct Exchange4.2 Topic Exchange4.3 Fanout Exchange5. Message 消息1. 初识RabbitMQ RabbitMQ 是一个开源的消息代理和队列服务器&#xff0c;用来通过普通协议在完全不同的应用之间共享数据&a…

接收并解析消息体传参、解析 json 参数

前些天发现了一个巨牛的人工智能学习网站&#xff0c;通俗易懂&#xff0c;风趣幽默&#xff0c;忍不住分享一下给大家。点击跳转到教程。 1.场景&#xff1a;postman 发送了一个 post 请求&#xff0c;如下&#xff1a; 2. 解析方式为用一个 vo 对象来接收 json。把 json 中的…

OpenCL memory object 之 传输优化

首先我们了解一些优化时候的术语及其定义&#xff1a; 1、deferred allocation&#xff08;延迟分配&#xff09;&#xff0c; 在第一次使用memory object传输数据时&#xff0c;runtime才对memory object真正分配空间。 这样减少了资源浪费&#xff0c;但第一次使用时要慢一些…