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();
if(y>=(imgHeight))
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}; const int s = ;
const int nStart = ; float lines[];
for(int i=;i<;i++)
lines[i] = read_imagef( source, sampler, (int2) (i-, y) ).x; for(int j=;j<imgWidth;){
float sum = lines[nStart] * m_nFilter[nStart];
#define GaussianTwoLines(m) \
sum += ( (lines[m] + lines[s--m])*m_nFilter[m] );
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines() write_imagef( dest, (int2) (j, y), sum ); for(int i = ; i<s-; i++) lines[i] = lines[i+];
j++;
lines[s-] = read_imagef( source, sampler, (int2) (j+, 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();
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}; const int s = ;
const int nStart = ; float lines[];
for(int i=;i<;i++)
lines[i] = read_imagef( source, sampler, (int2) (x ,i-) ).x; for(int j=;j<imgHeight;){
float sum = lines[nStart] * m_nFilter[nStart];
#define GaussianTwoLines(m) \
sum += ( (lines[m] + lines[s--m])*m_nFilter[m] );
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines()
GaussianTwoLines() write_imagef( dest, (int2) (x, j), sum ); for(int i = ; i<s-; i++) lines[i] = lines[i+];
j++;
lines[s-] = read_imagef( source, sampler, (int2) (x,j+) ).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(,,)))
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();
if(y>=(imgHeight))
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 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+,y),r(x+,y),r(x+,y),r(x+,y),r(x+,y),r(x+,y),r(x+,y),r(x+,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+, 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 line0 = r16(-,y);
for(int j=;j<imgWidth;){
float16 line1 = r16(j-+,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[];
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;
w16(j,y,sum );
j+=;
} } __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 line0 = rv16(x,-);
for(int j=;j<imgHeight;){
float16 line1 = rv16(x,j-+); 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[];
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;
wv16(x,j,sum );
j+=;
}
}