opencl gauss filter优化(一)

最新推荐文章于 2024-10-09 20:21:37 发布
最新推荐文章于 2024-10-09 20:21:37 发布
阅读量408 收藏
点赞数
原文链接：http://www.cnblogs.com/mlj318/p/5039348.html
版权

Platform: LG G3, Adreno 330 ,img size 3264x2448


C code

neon

GPU

300

60

29

 单位:ms


1. 目前按如下行列分解的方式最快29ms,Horizontal kernel globalWorksize[1] = {height+256-height%256};Vertical kernel globalWorksize2[1] = {width+256-width%256};

localWorksize2[] = {64}; localWorksize2 手动设为64时最快。

Porfile的结果为:Horizontal kernel 的wait time 有11ms,实际rum time 18ms.

这个wait time是什么呢？注释掉Horizontal kernel中的 vstore16(convert_uchar16(sum>>(ushort)8),0,pOutLine+j) ; 则wait time只有0.x ms.并且 localWorksize 越小wait time越长,为1时达到200ms,16时20ms. 难道是写内存等待时间，没有足够的ALU指令隐藏访存延时？写内存后进入下一个for循环，马上又读内存，所以没有ALU指令隐藏这个延时。然而Horizontal kernel的profile结果实际run time只有0.x ms,所有时间基本都是在wait.(更正：注释掉vstore16后,sum的计算被优化掉了，0.x ms是读内存的时间)

 

__kernel void ImageGaussianFilterHorizontal(__global const uchar* restrict source, // Source image
                            __global uchar* restrict  dest,  // Intermediate dest image
                                             const int imgWidth ,                // Image width
                                             const int imgHeight)
{
    const int y = get_global_id(0);
    if(y>=(imgHeight))
        return;
    const uchar m_nRightShiftNum = 8;
    const uchar Rounding = (1 << (m_nRightShiftNum - 1));
    const uchar  m_nFilter[11] = {1,4,8,16,32,134,32,16,8,4,1};

    const int s = 11;
    const int nStart = 5;
    const int nWidth = imgWidth;

    __global const uchar* pInLine = source + y*nWidth;
    __global uchar* pOutLine = dest + y*nWidth;
    int j;
    for(j = 0; j < nStart; j ++)
    {
        ushort sum = 0;

        for (int m = 0; m<s / 2; m++)
        {
            int k1 = (j + m - nStart);
            k1 = k1<0 ? -k1 : k1;

            int k2 = (j + nStart - m );
            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];
        }
        sum += pInLine[j] * m_nFilter[s / 2];
        sum = (sum + Rounding) >> 8;
        pOutLine[j] = (uchar)clamp(sum,(ushort)0,(ushort)255);
    }

    for ( ; (j+16)<= (nWidth - nStart); j+=16)
    {        
#define GAUSSIAN_LINE_NEON(m) \
sum += ( convert_ushort16(vload16(0,pInLine+j-nStart+m))* m_nFilter[m] );

        ushort16 sum =  (convert_ushort16(vload16(0,pInLine+j-nStart)) * m_nFilter[0]);
        GAUSSIAN_LINE_NEON(1);
        GAUSSIAN_LINE_NEON(2);
        GAUSSIAN_LINE_NEON(3);
        GAUSSIAN_LINE_NEON(4);
        GAUSSIAN_LINE_NEON(5);
        GAUSSIAN_LINE_NEON(6);
        GAUSSIAN_LINE_NEON(7);
        GAUSSIAN_LINE_NEON(8);
        GAUSSIAN_LINE_NEON(9);
        GAUSSIAN_LINE_NEON(10);

        sum += (ushort)Rounding;
        vstore16(convert_uchar16(sum>>(ushort)8),0,pOutLine+j) ;
    }

    for( ; j < nWidth; j ++)
    {
        ushort sum = 0;

        for (int m = 0; m<s / 2; m++)
        {
            int k1 = (j + m - nStart);

            int k2 = (j + nStart - m );
            k2 = k2 >= nWidth ? 2 * nWidth - 2 - k2 : k2;
            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];
        }
        sum += pInLine[j] * m_nFilter[s / 2];
        sum = (sum + Rounding) >> m_nRightShiftNum;
        pOutLine[j] =  (uchar)clamp(sum,(ushort)0,(ushort)255);
    }
}


__kernel void ImageGaussianFilterVertical( __global uchar* restrict source,   // Intermediate image processed by ImageGaussianFilterHorizontal()
                        __global uchar* restrict dest,  // Final destination image
                        const int imgWidth,
                                         const int imgHeight
                                    )
{
    const int x = get_global_id(0);
    if(x>=(imgWidth))
        return;
    const int x_offset = x;

    const int s = 11;
    const int nStart = s / 2;
    const int m_nRightShiftNum = 8;
    const int Rounding = (1 << (m_nRightShiftNum - 1));
    const uchar  m_nFilter[11] = {1,4,8,16,32,134,32,16,8,4,1};

    int y;
//    mem_fence(CLK_LOCAL_MEM_FENCE);

    ushort lines[11];
    lines[nStart] = (ushort)( source[x_offset]  );
    for(y=1;y<=nStart;y++)
    {
        lines[nStart+y] = (ushort)( source[y*imgWidth+x_offset]  );
        lines[nStart-y] = lines[nStart+y];
    }

    for(y=0;y<(imgHeight-nStart-1);)
    {

        ushort sum = lines[nStart] * m_nFilter[nStart];
#define    GaussianTwoLines(m) \
    sum += ( (lines[m] + lines[s-1-m])*m_nFilter[m] );
        GaussianTwoLines(0)
        GaussianTwoLines(1)
        GaussianTwoLines(2)
        GaussianTwoLines(3)
        GaussianTwoLines(4)

        sum += (ushort)Rounding;
        dest[y*imgWidth+x_offset]  = (uchar)(sum>>(ushort)8);

        y++;
        for(int i = 0; i<s-1; i++) lines[i] = lines[i+1];
        
        lines[s-1] =  (ushort)( source[(y+nStart)*imgWidth+x_offset]  );
        
    }

    for(y=imgHeight-nStart-1;y<(imgHeight-1);)
    {
        ushort sum = lines[nStart] * m_nFilter[nStart];
        GaussianTwoLines(0)
        GaussianTwoLines(1)
        GaussianTwoLines(2)
        GaussianTwoLines(3)
        GaussianTwoLines(4)

        sum += (ushort)Rounding;
        dest[y*imgWidth+x_offset]  = (uchar)(sum>>(ushort)8);

        y++;
        for(int i = 0; i<s-1; i++) {
            lines[i] = lines[i+1];
        }
        lines[s-1] = lines[(imgHeight-y)*2-2] ; //
    }
    //last y=imgHeight-1
    ushort sum = lines[nStart] * m_nFilter[nStart];
    GaussianTwoLines(0)
    GaussianTwoLines(1)
    GaussianTwoLines(2)
    GaussianTwoLines(3)
    GaussianTwoLines(4)

    sum += (ushort)Rounding;
    dest[y*imgWidth+x_offset]  = (uchar)(sum>>(ushort)8);
}
kernel

 

2.Horizontal kernel改进，预先load 2x16个所需的pixel,计算时从中提取,这样每次循环只需读一次内存。需要26ms,wait time 8ms.

 

    ushort16 line0 =  convert_ushort16(vload16(0,pInLine+j-nStart));
    for ( ; (j+16)<= (nWidth - nStart); j+=16)
    {
        ushort16 line1 =  convert_ushort16(vload16(0,pInLine+j-nStart+16));

        ushort16 temp0;
        ushort16 temp1;
        temp0 = line0;
        temp1.s0123 = line0.sabcd;
        temp1.s45 = line0.sef;
        temp1.s67 = line1.s01;
        temp1.s89abcdef = line1.s23456789;
        ushort16 sum =  ( temp0 + temp1 ) * m_nFilter[0];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s0;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s9;
        sum += ( temp0 +  temp1 ) * m_nFilter[1];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s1;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s8;
        sum += ( temp0 +  temp1 ) * m_nFilter[2];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s2;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s7;
        sum += ( temp0 +  temp1 ) * m_nFilter[3];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s3;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s6;
        sum += ( temp0 +  temp1 ) * m_nFilter[4];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s4;
        sum += ( temp0 ) * m_nFilter[5];

        sum += (ushort)Rounding;
        line0 = line1;
        vstore16(convert_uchar16(sum>>(ushort)8),0,pOutLine+j) ;
    }
View Code

 

3.不计算，只读写内存测试。那么wait time 3.2 ms,run time 18.2 ms.说明Horizontal kernel 耗时的极限也需3.2ms. 但是只是注释掉vstore16，还保留了读和计算，反而wait time还只有0.x ms，这又是为何？是读几乎没有wait,3.2ms都是写的wait time? (更正：注释掉vstore16后,sum的计算被优化掉了，0.x ms是读内存的时间)

a.再次测试，只有读wait time 0.xms ,只有写wait time 3.2ms.写比读的周期长.

for ( ; (j+16)<= (nWidth - nStart); j+=16)

{

ushort16 line1 = convert_ushort16(vload16(0,pInLine+j-nStart+16));

vstore16(0,0,pOutLine+j) ;

}

b.另外发现使用*((__global uint4*)(pOutLine+j)) = as_uint4(result);比vstore16快，wait time 2.5ms.高通 80-N8592-1_L_OpenCL_Programming_Guide 中提到：

Vectorized load/store of a larger data type is more optimal than a small data type; e.g., a load of uint2* is more optimal than uchar8* .

For optimal SP to L2 bandwidth performance, align read access to a 32-bit address and write access to a 128-bit address.

c.原来写的内存没有对齐，使用*((__global uint4*)(pOutLine+j-5)) = as_uint4(result);wait time 1.9ms.

d.最后加上sum计算，采用的Horizontal kernel如下，localWorksize[] = {64};时时间最少，需要23ms,wait time 4.7ms , localWorksize = 128时,wait 6ms.

并且使用__attribute__((work_group_size_hint(64,1,1))) ,耗时22ms.

 

__kernel __attribute__((work_group_size_hint(64,1,1)))   
void ImageGaussianFilterHorizontal(__global const uchar* restrict source, // Source image
                        __global uchar* restrict  dest,  // Intermediate dest image
                                             const int imgWidth ,                // Image width
                                             const int imgHeight)
{
    const int y = get_global_id(0);
    if(y>=(imgHeight))
        return;
    const uchar m_nRightShiftNum = 8;
    const uchar Rounding = (1 << (m_nRightShiftNum - 1));
    const uchar  m_nFilter[11] = {1,4,8,16,32,134,32,16,8,4,1};

    const int s = 11;
    const int nStart = 5;
    const int nWidth = imgWidth;


    __global const uchar* pInLine = source + y*nWidth;
    __global uchar* pOutLine = dest + y*nWidth;

    int j;
    uchar temp[5];
    for(j = 0; j < nStart; j ++)
    {
        ushort sum = 0;

        for (int m = 0; m<s / 2; m++)
        {
            int k1 = (j + m - nStart);
            k1 = k1<0 ? -k1 : k1;

            int k2 = (j + nStart - m );
            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];
        }
        sum += pInLine[j] * m_nFilter[s / 2];
        sum = (sum + Rounding) >> 8;
        temp[j] = (uchar)clamp(sum,(ushort)0,(ushort)255);
    }

    uchar16 result,pre_result;
    pre_result.sbcde = (uchar4)(temp[0],temp[1],temp[2],temp[3]);
    pre_result.sf = temp[4];

    ushort16 line0 =  convert_ushort16(vload16(0,pInLine+j-nStart));
    for ( ; (j+16)<= (nWidth - nStart); j+=16)
    {
        //prefetch(pInLine+j-nStart,32); //无变化
        ushort16 line1 =  convert_ushort16(vload16(0,pInLine+j-nStart+16));

        ushort16 temp0;
        ushort16 temp1;
        temp0 = line0;
        temp1.s0123 = line0.sabcd;
        temp1.s45 = line0.sef;
        temp1.s67 = line1.s01;
        temp1.s89abcdef = line1.s23456789;
        ushort16 sum =  ( temp0 + temp1 ) * m_nFilter[0];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s0;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s9;
        sum += ( temp0 +  temp1 ) * m_nFilter[1];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s1;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s8;
        sum += ( temp0 +  temp1 ) * m_nFilter[2];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s2;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s7;
        sum += ( temp0 +  temp1 ) * m_nFilter[3];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s3;
        temp1.s0123456789abcdef = temp1.s00123456789abcde;
        temp1.s0 = line0.s6;
        sum += ( temp0 +  temp1 ) * m_nFilter[4];
        temp0.s0123456789abcdef = temp0.s123456789abcdeff;
        temp0.sf = line1.s4;
        sum += ( temp0 ) * m_nFilter[5];

        sum += (ushort)Rounding;
        line0 = line1;

        result.s0123 = pre_result.sbcde;
        result.s4 = pre_result.sf;
        pre_result = convert_uchar16(sum>>(ushort)8) ;

        result.s5 = pre_result.s0;
        result.s67 = pre_result.s12;
        result.s89abcdef = pre_result.s3456789a;
        *( (__global uint4*)(pOutLine+j-5) ) =  (as_uint4)(result) ;
    }

    *( (__global uint*)(pOutLine+j-5) ) = (as_uint)(pre_result.sbcde);//last 5 bytes
    pOutLine[j-1] = pre_result.sf;

    for( ; j < nWidth; j ++)
    {
        ushort sum = 0;

        for (int m = 0; m<s / 2; m++)
        {
            int k1 = (j + m - nStart);

            int k2 = (j + nStart - m );
            k2 = k2 >= nWidth ? 2 * nWidth - 2 - k2 : k2;
            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];
        }
        sum += pInLine[j] * m_nFilter[s / 2];
        sum = (sum + Rounding) >> m_nRightShiftNum;
        pOutLine[j] =  (uchar)clamp(sum,(ushort)0,(ushort)255);
    }
}
View Code

 

 

 

 

 

转载于:https://www.cnblogs.com/mlj318/p/5039348.html

weixin_30376083
关注
性能优化-OpenCL kernel 开发
发狂的小花的博客
01-22 1839
OpenCL kernel 部分的介绍。 OpenCL kernel 是运行在设备端的，采用OpenCL C 语言进行开发，本文接下来首先给出一个简单的OpenCL kernel 样例，然后对OpenCL C 语言的各个部分做详细的说明，最后会给出一个完整的OpenCL程序实例，相信通过本文的学习之后大家应该可以在实际工作中使用OpenCL来优化程序的性能。
性能优化-OpenCL 介绍
发狂的小花的博客
01-22 1588
如今的时代是一个数据爆炸的时代，如何处理海量的数据是各行各业不得不面临的一个重要问题。随着摩尔定律的失效，支持大量并行计算的异构处理器开始大放异彩。其中，无论是在 PC 端还是移动端，GPU 无疑是其中最璀璨的明星。接下来，我们首先会对 GPU 的发展历史和硬件架构（移动端）进行简单地介绍，随后对 OpenCL 编程框架进行详细的介绍。
Android平台下OpenCL并行Sobel滤波实现
06-11
该压缩包为Eclipse工程文件，包含程序运行所需要的OpenCV apk文件。程序主要是利用OpenCL在智能手机（安卓系统）实现并行化，达到缩短处理时间、降低功耗的目的。
OpenCL水平垂直高斯滤波
weixin_45258796的博客
03-26 428
水平高斯滤波 垂直高斯滤波
openGauss开源数据库实战二
最新发布
weixin_73321443的博客
10-09 1017
在CentOS单机上安装openGauss DBMS，为学习openGauss DBMS准备数据库环境。
opencl _高斯核卷积
beibei_hao_csdn的博客
01-17 805
生成高斯核代码： void GetGaussianKernel(float *gaus, const int size,const float sigma)  {    const double PI=4.0*atan(1.0); //圆周率π赋值    int center=size/2;    double sum=0;    for(int i=0;i   { 
【高性能计算】opencl语法及相关概念（四）:结合opencv进行图像高斯模糊处理
种一棵树最好的时间是十年前，其次是现在
08-31 938
高斯模糊是一种常用的图像处理技术，用于减少图像中的噪点和细节，并实现图像的平滑效果。它是基于高斯函数的卷积操作，通过对每个像素周围的邻域像素进行加权平均来实现模糊效果。具体而言，高斯模糊通过在图像上滑动一个卷积核，将卷积核与输入图像的对应像素进行一一相乘，并将结果相加，从而产生输出图像的每个像素值。这个卷积核是一个二维高斯函数，它的形状决定了模糊的程度。在高斯函数中，离中心像素越远的像素会被赋予更小的权重，从而降低了离中心像素的贡献，实现模糊的效果。通过调整高斯核的大小和标准差参数，可以控制模糊的程度。
学习OpenCL与卷积神经网络 Day5——基于OpenCL的图像旋转以及高斯过滤的程序构建
不偷懒会死星人的blog
12-17 914
目录前言采样器设备端采样器创建主机端采样器创建图像旋转内核创建主机函数创建完整程序旋转输出结果高斯过滤内核创建主机函数创建完整程序高斯模糊输出结果 前言 OpenCL其实算是一种“加速语言”，其由设备上运行的kernel函数语言和控制平台的API组成，它通过将某些简单而又重复的工作转交给GPU/FPGA外接设备，实现异构并行来加速原本的工作。比如在OpenCV中就引入了OCL module，其编译模块要比CUDA快很多，它就起到了一个加速图像处理过程的作用。当然，今天的工作只是简单的熟悉一下如何在OpenC
Opencl 之bilateral filter优化
Emel 的博客
07-19 1437
由于项目的需要在我还没有搞清楚Opencl框架的情况下就开始去用opencl方式去优化一个图像处理算法--bilateral_filter，可想而知其过程的艰辛，好在昨天已经搞定收工了，感谢这一个月的过程让我从一个opencl路人变成了亲戚，特此写一篇博客记录下一些心得。废话少说现在开始。 拿到这个任务是我首先感觉相对我之前写代码还是比较高大上的，因此心里也暗暗下决心搞定这个问题，于是先去了解双...
OpenCL编程与优化Part II
05-29
标题和描述中提到的知识点主要涵盖了OpenCL编程和优化方面的内容，具体可以总结为以下几个方面： 1. OpenCL在GPU上的性能瓶颈 OpenCL编程在CPU和GPU协同工作的环境中，存在一些性能瓶颈。其中最明显的是CPU和GPU...
性能优化-一文宏观理解OpenCL
发狂的小花的博客
01-11 1039
OpenCL概述 OpenCL架构与原理 OpenCL开发环境搭建 OpenCL数据类型与运算 OpenCL并行计算优化 OpenCL图像处理应用 OpenCL性能分析与调试
基于CUDA的并行计算技术+opencv完成图像高斯滤波和双边滤波
05-22
基于CUDA的并行计算技术+opencv完成图像高斯滤波和双边滤波，开发版本为VS2019+openCV3.4
理想，高斯滤波器
04-23
这是一个基于MATLAB的，理想，高斯，滤波器程序
opencl实现矩阵加法
11-06
opencl实现矩阵加法，非常适合初学者学习使用，能够全面了解OpenCL的运行原理，本实验环境ubuntu+NVIDIA CUDA8.0，关于环境配置可以查看我的博文。
OpenCLDemo
05-30
一个OpenCL小程序，适合OpenCL初学者学习使用。 a simple program for study.
opencl高斯源码整理
jacke121的专栏
03-18 2559
opencl高斯源码整理
OpenCL双边滤波实现美颜功能
weixin_30426957的博客
02-22 299
OpenCL是一个并行异构计算的框架，包括intel，AMD，英伟达等等许多厂家都有对它的支持，不过英伟达只到1.2版本，主要发展自己的CUDA去了。虽然没有用过CUDA，但个人感觉CUDA比OpenCL更好一点，但OpenCL支持面更管，CPU,GPU,DSP,FPGA等多种芯片都能支持OpenCL。OpenCL与D3D中的像素着色器非常相似。 1.双边滤波原理 双边滤波器的...
引导滤波guideFilter原理推导与实验
thequitesunshine007的博客
01-25 481
转载至： https://zhuanlan.zhihu.com/p/98368439 https://zhuanlan.zhihu.com/p/161666126 论文原文 Kaiming He, Jian Sun, Xiaoou Tang. Guided Image Filtering[J]. IEEE Transactions on Pattern Analysis & Machine Intelligence, 2013,35(6):1397-1409. 导向滤波（Guided Flite
OpenCL 数据传输操作
INFINALGEORGE的博客
03-30 1703
1.加载和保存同类型的数据:如果发送端和接收端的数据类型相同，例如将一个区域的int4型向量加载到另一区域中，只知道“=”即可。除了对变量赋值外，等号还可以将一个区域的数据赋值传输到另一个区域中。例如，内核想将全局内存中的数据加载到局部内存中，并将处理完的结果保存回全局内存。如果要将局部内存改为私有内存，只需要将数据声明中的_local标识符去掉即可，默认的标识符便是_private,因此，如果没...
OpenCL模型的GPU性能优化技术研究
"面向OpenCL模型的GPU性能优化" OpenCL（Open Computing Language）是一种开放标准，用于编写跨平台的并行代码，特别是在GPU（图形处理器）上进行高性能计算。本文主要探讨了如何针对OpenCL模型优化GPU的性能，以...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值