1.根据前两次的最终结果：

使用普通buffer,Horizontal 5ms, Vertical 17 ms

使用image buffer：Horizontal 9.4ms, Vertical 6.4 ms

那么使用 Horizontal普通buffer,Vertical image buffer 组合方式的话，是不是时间最少？只是Intermediate image仍使用image对象，Horizontal kernel中的写操作需要改变。

结果： Horizontal 的最大local_work_size只能是32, Horizontal 增至8ms, Vertical 6.4ms

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

#define r(xc,y) read_imagef( source, sampler,  (int2) (xc, y) ).x

#define w16(x,y,sum) write_imagef( dest, (int2) (x, y), sum.s0 );write_imagef( dest, (int2) (x+1, y), sum.s1 );\

        write_imagef( dest, (int2) (x+, y), sum.s2 );write_imagef( dest, (int2) (x+, y), sum.s3 );\

        write_imagef( dest, (int2) (x+, y), sum.s4 );write_imagef( dest, (int2) (x+, y), sum.s5 );\

        write_imagef( dest, (int2) (x+, y), sum.s6 );write_imagef( dest, (int2) (x+, y), sum.s7 );\

        write_imagef( dest, (int2) (x+, y), sum.s8 );write_imagef( dest, (int2) (x+, y), sum.s9 );\

        write_imagef( dest, (int2) (x+, y), sum.sa );write_imagef( dest, (int2) (x+, y), sum.sb );\

        write_imagef( dest, (int2) (x+, y), sum.sc );write_imagef( dest, (int2) (x+, y), sum.sd );\

        write_imagef( dest, (int2) (x+, y), sum.se );write_imagef( dest, (int2) (x+, y), sum.sf );

__kernel __attribute__((work_group_size_hint(,,)))

void ImageGaussianFilterHorizontal(__global const uchar* restrict source, // Source image

                                    __write_only image2d_t   dest,  // Intermediate dest image

                                     const int imgWidth ,                // Image width

                                     const int imgHeight)

{

    const int y = get_global_id();

    if(y>=(imgHeight))

        return;

    const uchar m_nRightShiftNum = ;

    const uchar Rounding = ( << (m_nRightShiftNum - ));

    const uchar  m_nFilter[] = {,,,,,,,,,,};

    const int s = ;

    const int nStart = ;

    const int nWidth = imgWidth;

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

    int j;

    for(j = ; j < nStart; j ++)

    {

        ushort sum = ;

        for (int m = ; m<s / ; m++)

        {

            int k1 = (j + m - nStart);

            k1 = k1< ? -k1 : k1;

            int k2 = (j + nStart - m );

            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];

        }

        sum += pInLine[j] * m_nFilter[s / ];

        //sum = (sum + Rounding) >> 8;

        write_imagef( dest, (int2) (j, y), convert_float(sum)/(255.0*) );

    }

    ushort16 line0 =  convert_ushort16(vload16(,pInLine+j-nStart));

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

    {

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

        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[];

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;

        temp0.sf = line1.s0;

        temp1.s0123456789abcdef = temp1.s00123456789abcde;

        temp1.s0 = line0.s9;

        sum += ( temp0 +  temp1 ) * m_nFilter[];

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;

        temp0.sf = line1.s1;

        temp1.s0123456789abcdef = temp1.s00123456789abcde;

        temp1.s0 = line0.s8;

        sum += ( temp0 +  temp1 ) * m_nFilter[];

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;

        temp0.sf = line1.s2;

        temp1.s0123456789abcdef = temp1.s00123456789abcde;

        temp1.s0 = line0.s7;

        sum += ( temp0 +  temp1 ) * m_nFilter[];

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;

        temp0.sf = line1.s3;

        temp1.s0123456789abcdef = temp1.s00123456789abcde;

        temp1.s0 = line0.s6;

        sum += ( temp0 +  temp1 ) * m_nFilter[];

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;

        temp0.sf = line1.s4;

        sum += ( temp0 ) * m_nFilter[];

        line0 = line1;

        float16 sum2 = (convert_float16(sum))/(255.0*);

        w16(j,y,sum2 );

    }

    for( ; j < nWidth; j ++)

    {

        ushort sum = ;

        for (int m = ; m<s / ; m++)

        {

            int k1 = (j + m - nStart);

            int k2 = (j + nStart - m );

            k2 = k2 >= nWidth ?  * nWidth -  - k2 : k2;

            sum += (pInLine[k1] + pInLine[k2])*m_nFilter[m];

        }

        sum += pInLine[j] * m_nFilter[s / ];

        //sum = (sum + Rounding) >> m_nRightShiftNum;

        write_imagef( dest, (int2) (j, y), convert_float(sum)/(255.0*) );

    }

}

2.使用各种办法,最终也只能降到13.7ms,Horizontal 7.5, Vertical 6ms,最终代码如下.

更新：H和V都 去掉__attribute__ 属性,local_work_size都设置NULL,让opencl自己选择，H 的最大local_work_size又变回了64，总时间13ms.因为在LG G4,adreno 418上运行却需要40ms,在adreno 418上的local_work_size最大可以是1024，却被强制设成了32.

a.使用mad指令做sum乘加，结果有误差，时间也略增.fma 是无限精度,mad 是快速方法，结果是近似值。

b.使用 pInTemp 读fisrt 16 bytes,避免重复读取，有0.x ms的优势

c.边界使用了mirror repeat

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

#define r(xc,y) read_imagef( source, sampler,  (int2) (xc, y) ).x

#define w16(x,y,sum) write_imagef( dest, (int2) (x, y), sum.s0 );write_imagef( dest, (int2) (x+1, y), sum.s1 );\

        write_imagef( dest, (int2) (x+, y), sum.s2 );write_imagef( dest, (int2) (x+, y), sum.s3 );\

        write_imagef( dest, (int2) (x+, y), sum.s4 );write_imagef( dest, (int2) (x+, y), sum.s5 );\

        write_imagef( dest, (int2) (x+, y), sum.s6 );write_imagef( dest, (int2) (x+, y), sum.s7 );\

        write_imagef( dest, (int2) (x+, y), sum.s8 );write_imagef( dest, (int2) (x+, y), sum.s9 );\

        write_imagef( dest, (int2) (x+, y), sum.sa );write_imagef( dest, (int2) (x+, y), sum.sb );\

        write_imagef( dest, (int2) (x+, y), sum.sc );write_imagef( dest, (int2) (x+, y), sum.sd );\

        write_imagef( dest, (int2) (x+, y), sum.se );write_imagef( dest, (int2) (x+, y), sum.sf );

//line0 start from j-5,line1 from j-5+16

#define GaussianShift16 {\

        temp0 = line0;\

        temp1.s0123 = line0.sabcd;\

        temp1.s45 = line0.sef;\

        temp1.s67 = line1.s01;\

        temp1.s89abcdef = line1.s23456789;\

        sum =  ( temp0 + temp1 ) * m_nFilter[];\

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;\

        temp0.sf = line1.s0;\

        temp1.s0123456789abcdef = temp1.s00123456789abcde;\

        temp1.s0 = line0.s9;\

        sum += ( temp0 +  temp1 ) * m_nFilter[];\

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;\

        temp0.sf = line1.s1;\

        temp1.s0123456789abcdef = temp1.s00123456789abcde;\

        temp1.s0 = line0.s8;\

        sum += ( temp0 +  temp1 ) * m_nFilter[];\

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;\

        temp0.sf = line1.s2;\

        temp1.s0123456789abcdef = temp1.s00123456789abcde;\

        temp1.s0 = line0.s7;\

        sum += ( temp0 +  temp1 ) * m_nFilter[];\

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;\

        temp0.sf = line1.s3;\

        temp1.s0123456789abcdef = temp1.s00123456789abcde;\

        temp1.s0 = line0.s6;\

        sum += ( temp0 +  temp1 ) * m_nFilter[];\

        temp0.s0123456789abcdef = temp0.s123456789abcdeff;\

        temp0.sf = line1.s4;\

        sum += ( temp0 ) * m_nFilter[];}

__kernel __attribute__((work_group_size_hint(,,)))

void ImageGaussianFilterHorizontal(__global const uchar* restrict source, // Source image

                                    __write_only image2d_t   dest,  // Intermediate dest image

                                     const int imgWidth ,                // Image width

                                     const int imgHeight)

{

    const int y = get_global_id();

    if(y>=(imgHeight))

        return;

    const uchar  m_nFilter[] = {,,,,,,,,,,};

    const int s = ;

    const int nStart = ;

    __global const uchar* pInLine = source + y*imgWidth;

    int j;

    uchar pInTemp[];

    *( (uint4*)(pInTemp) ) = *((__global uint4*)(pInLine)) ;//first 16 bytes

    for(j = ; j < nStart; j ++)

    {

        ushort sum = ;

        for (int m = ; m<s / ; m++)

        {

            int k1 = (j + m - nStart);

            k1 = k1< ? -k1 : k1;

            int k2 = (j + nStart - m );

            sum += (pInTemp[k1] + pInTemp[k2])*m_nFilter[m];

        }

        sum += pInTemp[j] * m_nFilter[s / ];

        write_imagef( dest, (int2) (j, y), convert_float(sum)/(255.0*) );

    }

    ushort16 temp0;

    ushort16 temp1;

    ushort16 sum;

    ushort16 line0,line1;

    line0 =  convert_ushort16(*((uchar16*)pInTemp));

    for ( ; j< (imgWidth-); j+=)

    {

        line1 = convert_ushort16(vload16(,pInLine+j-nStart+));//convert_ushort16( as_uchar16(*((__global uint4*)(pInLine+j-nStart+16))) ) ;

        GaussianShift16

        line0 = line1;

        float16 sum2 = (convert_float16(sum))/(255.0*);

        w16(j,y,sum2 );

    }

    {

        //last 16 pixel,some pixels may caculate again

        j = imgWidth-;

        line0 =  convert_ushort16(vload16(,pInLine+j-nStart));

        //mirror repeat read

        line1.s0123 =  convert_ushort4( vload4(,pInLine+imgWidth-nStart) );

        line1.s4567 = (ushort4)( pInLine[imgWidth-],line1.s3,line1.s21 ) ;

        line1.s89 = (ushort2)(line1.s0,line0.sf);

        GaussianShift16

        float16 sum2 = (convert_float16(sum))/(255.0*);

        w16(j,y,sum2 );

    }

}

__kernel  __attribute__((work_group_size_hint(,,)))

void ImageGaussianFilterVertical(__read_only image2d_t  source, // Source image

                                __write_only image2d_t   dest,

                                 const int imgWidth ,

                                 const int imgHeight)

{

    const int x = get_global_id();

    if(x>=(imgWidth))

        return;

    const float m_nFilter[] = {/256.0,/256.0,/256.0,/256.0,/256.0,/256.0,/256.0,/256.0,/256.0,/256.0,/256.0};

#define rv16(x,y) (float16)( r(x,y),r(x,y+1),r(x,y+2),r(x,y+3),r(x,y+4),r(x,y+5),r(x,y+6),r(x,y+7),\

                r(x,y+),r(x,y+),r(x,y+),r(x,y+),r(x,y+),r(x,y+),r(x,y+),r(x,y+))

#define wv16(x,y,sum) write_imagef( dest, (int2) (x,y), sum.s0 );write_imagef( dest, (int2) (x,y+1), sum.s1 );\

        write_imagef( dest, (int2) (x,y+), sum.s2 );write_imagef( dest, (int2) (x,y+), sum.s3 );\

        write_imagef( dest, (int2) (x,y+), sum.s4 );write_imagef( dest, (int2) (x,y+), sum.s5 );\

        write_imagef( dest, (int2) (x,y+), sum.s6 );write_imagef( dest, (int2) (x,y+), sum.s7 );\

        write_imagef( dest, (int2) (x,y+), sum.s8 );write_imagef( dest, (int2) (x,y+), sum.s9 );\

        write_imagef( dest, (int2) (x,y+), sum.sa );write_imagef( dest, (int2) (x,y+), sum.sb );\

        write_imagef( dest, (int2) (x,y+), sum.sc );write_imagef( dest, (int2) (x,y+), sum.sd );\

        write_imagef( dest, (int2) (x,y+), sum.se );write_imagef( dest, (int2) (x,y+), sum.sf );

    float16 temp0;

    float16 temp1;

    float16 sum;

    float16 line0,line1;

    line0 =  rv16(x,-);

    line0.s0123 = line0.sa987;//mirror repeat

    line0.s4 = line0.s6;

    int j;

    for(j=;j<imgHeight-;j+=){

        line1 =  rv16(x,j-+);

        GaussianShift16

        line0 = line1;

        wv16(x,j,sum );

    }

    //last 16 pixel,some pixels may caculate again if imgHeight not 16 bytes align

    j = imgHeight-;

    line0 =  rv16(x,j-);

    //mirror repeat read

    const int y = imgHeight-;

    line1.s0123 =  (float4)( r(x,y),r(x,y+),r(x,y+),r(x,y+) );

    line1.s4567 = (float4)( r(x,y+),line1.s3,line1.s21 );

    line1.s89 = (float2)(line1.s0,line0.sf);

    GaussianShift16

    wv16(x,j,sum );

}

总结：1.local_work_size 对时间的影响比较大,有时使用NULL默认的就可以，有时需要一个个去试。

使用vector 类型,local memory,kernel代码结构 都会对 local_work_size 最大值有影响

2.profile中的wait time可能是读写memory还有其它的等待时间，rum time是ALU计算执行的时间。

3.避免对global memory的重复读写,预先缓存下来再用

4.image buffer的读写比普通buffer快，也没有按行按列读写的效率差异.尽量使用image buffer

5.read/write_imageui 并不比 read/write_imagef 快，一般就使用float

6.write 比read 要慢很多，内存未对齐也会慢些

7.使用vector 读写,计算 都会更快.image buffer虽然是单点读，组合成vector计算也更快.

8.half类型存在精度问题，会引入误差，在这里也不比float快

9.如果不确定local_work_size,就设置成NULL,让opencl自己选择。

不同的GPU 上local_work_size最大值不一样，比如这个kernel 在Adreno 330上最大64,在adreno 418上最大1024.

最新文章

  1. RN中listView的每个item等高
  2. 在web上逐行输出较大的txt文件
  3. Oracle转移数据表空间存储位置
  4. .Net 闭包理解
  5. ruby的在ubuntu上的安装
  6. mysql /*! 50100 ... */ 条件编译
  7. 粒子系统1：简介&amp;工具使用
  8. 写个简单的ANT脚本来编译项目
  9. python-Day4-迭代器-yield异步处理--装饰器--斐波那契--递归--二分算法--二维数组旋转90度--正则表达式
  10. js导出table到excel，同时兼容FF和IE
  11. 简单的jquery ajax文件上传功能
  12. WPS怎么让前几页的页眉或者页脚与后面的不同
  13. JAVA反射优化
  14. Ubuntu 16.04 启用 点击Launcher图标，窗口实现最小化 功能
  15. OO第一单元表达式求导作业总结
  16. 5DAY高级权限
  17. URL记录
  18. [学习笔记]凸优化/WQS二分/带权二分
  19. SQL Server - 文件组，文件，备份，分区
  20. [Java] 理解JVM之三：垃圾回收机制

热门文章

  1. 更新yum源
  2. Linux(centOS6.5)下SVN的安装、配置及开机启动
  3. 第十四章：高级I/O
  4. 客服端调用自定义宿主的WCF报错&quot;没有终结点在侦听可以接受消息的http://localhost:8085/mex。这通常是由于不正确的地址或者 SOAP 操作导致的错误&quot;的解决方案。
  5. log4j.properties详解与例子
  6. 第二课 android项目结构
  7. Python3基础 print 自带换行功能
  8. 高通平台msm8909 LK 实现LCD 兼容
  9. SQLServer使用表值参数，高性能批量插入数据
  10. 【转载】UML类图知识整理