opencl gauss filter优化(三)

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+2, y), sum.s2 );write_imagef( dest, (int2) (x+3, y), sum.s3 );\
        write_imagef( dest, (int2) (x+4, y), sum.s4 );write_imagef( dest, (int2) (x+5, y), sum.s5 );\
        write_imagef( dest, (int2) (x+6, y), sum.s6 );write_imagef( dest, (int2) (x+7, y), sum.s7 );\
        write_imagef( dest, (int2) (x+8, y), sum.s8 );write_imagef( dest, (int2) (x+9, y), sum.s9 );\
        write_imagef( dest, (int2) (x+10, y), sum.sa );write_imagef( dest, (int2) (x+11, y), sum.sb );\
        write_imagef( dest, (int2) (x+12, y), sum.sc );write_imagef( dest, (int2) (x+13, y), sum.sd );\
        write_imagef( dest, (int2) (x+14, y), sum.se );write_imagef( dest, (int2) (x+15, y), sum.sf );

__kernel __attribute__((work_group_size_hint(32,1,1)))
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(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;

    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;
        write_imagef( dest, (int2) (j, y), convert_float(sum)/(255.0*256) );
    }

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

        line0 = line1;

        float16 sum2 = (convert_float16(sum))/(255.0*256);
        w16(j,y,sum2 );
    }

    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;
        write_imagef( dest, (int2) (j, y), convert_float(sum)/(255.0*256) );
    }

}

View Code