opencl gauss filter优化(二)

1.buffer使用image的方式：Horizontal 与 Vertical 算法一样, 共需30ms,wait time 19ms.

 

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
__kernel void ImageGaussianFilterHorizontal(__read_only image2d_t  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 float m_nFilter[11] = {1/256.0,4/256.0,8/256.0,16/256.0,32/256.0,134/256.0,32/256.0,16/256.0,8/256.0,4/256.0,1/256.0};

    const int s = 11;
    const int nStart = 5;
    
    float lines[11];
    for(int i=0;i<11;i++)
        lines[i] = read_imagef( source, sampler,  (int2) (i-5, y) ).x;

    for(int j=0;j<imgWidth;){
    float 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)    
        
        write_imagef( dest, (int2) (j, y), sum );
        
        for(int i = 0; i<s-1; i++) lines[i] = lines[i+1];
        j++;
        lines[s-1] = read_imagef( source, sampler, (int2) (j+5, y) ).x;
    }
}

__kernel 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(0);
    if(x>=(imgWidth))
        return;
    const float m_nFilter[11] = {1/256.0,4/256.0,8/256.0,16/256.0,32/256.0,134/256.0,32/256.0,16/256.0,8/256.0,4/256.0,1/256.0};

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

    float lines[11];
    for(int i=0;i<11;i++)
        lines[i] = read_imagef( source, sampler,  (int2) (x ,i-5) ).x;

    for(int j=0;j<imgHeight;){
    float 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)

        write_imagef( dest, (int2) (x, j), sum );

        for(int i = 0; i<s-1; i++) lines[i] = lines[i+1];
        j++;
        lines[s-1] = read_imagef( source, sampler, (int2) (x,j+5) ).x;
    }
}

 

2.只运行 Horizontal 19ms,wait time 19ms. 注释掉 write_imagef 2.4ms(wait time,run time都是0.0xms)(更新：sum计算被优化,0.x ms就是读image的时间).

a.顺序调整为：

lines[s-1] = read_imagef( source, sampler, (int2) (j+5, y) ).x;

write_imagef( dest, (int2) (j-1, y), sum );

16.9ms,很奇怪sum用固定的0,0.2替代时间只有3.9ms?????把计算部分注释掉，只读写imgage,也是3.9ms, 计算sum的部分被编译器优化掉了？

b. if(sum>0)

lines[s-1] = read_imagef( source, sampler, (int2) (j+5, y) ).x;

write_imagef( dest, (int2) (j-1, y), 0.2 );

如此测试,17ms,看来是sum的计算被优化掉了.

c.if(sum>=0)

j++;

//lines[s-1] = read_imagef( source, sampler, (int2) (j+5, y) ).x;

//write_imagef( dest, (int2) (j-1, y), sum );

只计算,5.7ms,但还是wait time 5.7ms？？？

 

 

3.使用float16 vector 计算,总共耗时15.6 ms,wait time 9.3ms,rum time 6.3ms.使用 __attribute__ 能减少1ms以内.其中Horizontal:wait time 9.4ms,rum time 0.008ms ,Vertical:wait time 0.07ms,rum time 6.4ms.

 

不知道为什么使用fma指令替代sum+= ,需要近2s,而且localWorksize最大只能32.

 

使用half16 精度，反而还要17ms,而且结果有1-2的误差。

 

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

__kernel __attribute__((work_group_size_hint(64,1,1)))
void ImageGaussianFilterHorizontal(__read_only image2d_t  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 float m_nFilter[11] = {1/256.0,4/256.0,8/256.0,16/256.0,32/256.0,134/256.0,32/256.0,16/256.0,8/256.0,4/256.0,1/256.0};


#define r(xc,y) read_imagef( source, sampler,  (int2) (xc, y) ).x
#define r16(x,y) (float16)( r(x,y),r(x+1,y),r(x+2,y),r(x+3,y),r(x+4,y),r(x+5,y),r(x+6,y),r(x+7,y),
                r(x+8,y),r(x+9,y),r(x+10,y),r(x+11,y),r(x+12,y),r(x+13,y),r(x+14,y),r(x+15,y))

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

    float16 line0 =  r16(-5,y);
    for(int j=0;j<imgWidth;){
        float16 line1 =  r16(j-5+16,y);

        float16 temp0;
        float16 temp1;
        temp0 = line0;
        temp1.s0123 = line0.sabcd;
        temp1.s45 = line0.sef;
        temp1.s67 = line1.s01;
        temp1.s89abcdef = line1.s23456789;
        float16 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;
        w16(j,y,sum );
        j+=16;
    }

}

__kernel  __attribute__((work_group_size_hint(64,1,1)))
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(0);
    if(x>=(imgWidth))
        return;
    const float m_nFilter[11] = {1/256.0,4/256.0,8/256.0,16/256.0,32/256.0,134/256.0,32/256.0,16/256.0,8/256.0,4/256.0,1/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+8),r(x,y+9),r(x,y+10),r(x,y+11),r(x,y+12),r(x,y+13),r(x,y+14),r(x,y+15))

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

    float16 line0 =  rv16(x,-5);
    for(int j=0;j<imgHeight;){
        float16 line1 =  rv16(x,j-5+16);

        float16 temp0;
        float16 temp1;
        temp0 = line0;
        temp1.s0123 = line0.sabcd;
        temp1.s45 = line0.sef;
        temp1.s67 = line1.s01;
        temp1.s89abcdef = line1.s23456789;
        float16 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;
        wv16(x,j,sum );
        j+=16;
    }
}