0

我创建了一个无分支的中值 3x3 过滤器，并且在高分辨率图像（~4K x 3K）上每次通过大约 200 磨机，我认为如果我创建内核来利用工作组，我可以获得更好的时间. 不幸的是，工具告诉我我做错了什么，因为我想知道。

#define wgs 16
//Work group size
#define cas3(a, b)                              \
    do {                                            \
            float4 x = a;                           \
            int4 c = a> b;                          \
            a.s012 = select(b, a, c).s012;          \
            b.s012 = select(x, b, c).s012;          \
    } while (0)

__kernel void median3x3_rgb( read_only image2d_t src, write_only image2d_t dst) {
    int gx = get_global_id(0), gy = get_global_id(1);
    int lx = get_local_id(0), ly = get_local_id(1);
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE |    CLK_FILTER_NEAREST;

    if ((gx >= get_image_width(dst)) | (gy >= get_image_height(dst)))
            return;

__local float4 la[wgs+2][wgs+2];

la[lx+1][ly+1] = read_imagef(src,smp,(int2)(gx,gy));    

if(lx == 0){
    la[lx][ly+1] = read_imagef(src,smp,(int2)(gx-1,gy));
    if(ly == 0)
        la[lx+1][ly] = read_imagef(src,smp,(int2)(gx-1,gy));
    if(ly == wgs)
        la[lx+1][ly+2] = read_imagef(src,smp,(int2)(gx+1,gy));
}
else if(lx == wgs){
    la[lx+2][ly+1] = read_imagef(src,smp,(int2)(gx+1,gy));
    if(ly == 0)
        la[lx+1][ly] = read_imagef(src,smp,(int2)(gx-1,gy));
    if(ly == wgs)
        la[lx+1][ly+2] = read_imagef(src,smp,(int2)(gx+1,gy));
}
else if(ly == 0)
    la[lx+1][ly] = read_imagef(src,smp,(int2)(gx-1,gy));
else if(ly == wgs)
    la[lx+1][ly+2] = read_imagef(src,smp,(int2)(gx+1,gy));


barrier(CLK_LOCAL_MEM_FENCE); //----------------------- mem barrier

    float4 s0 = la[ lx-1][ly-1 ];
    float4 s1 = la[ lx  ][ly-1 ];
    float4 s2 = la[ lx+1][ly-1 ];
    float4 s3 = la[ lx-1][ly   ];
    float4 s4 = la[ lx  ][ly   ];
    float4 s5 = la[ lx+1][ly   ];
    float4 s6 = la[ lx-1][ly+1 ];
    float4 s7 = la[ lx  ][ly+1 ];
    float4 s8 = la[ lx+1][ly+1 ];

开始排序这个区域工作正常不用担心

    // stage0
    cas3(s1, s2);
    cas3(s4, s5);
    cas3(s7, s8);

    // 1
    cas3(s0, s1);
    cas3(s3, s4);
    cas3(s6, s7);

    // 2
    cas3(s1, s2);
    cas3(s4, s5);
    cas3(s7, s8);

    // 3/4
    cas3(s3, s6);
    cas3(s4, s7);
    cas3(s5, s8);
    cas3(s0, s3);

    cas3(s1, s4);
    cas3(s2, s5);
    cas3(s3, s6);

    cas3(s4, s7);
    cas3(s1, s3);

    cas3(s2, s6);
    cas3(s2, s3);
    cas3(s4, s6);

    cas3(s3, s4);

结束排序

    write_imagef(dst, (int2) (gx, gy), s4);
}
4

1 回答 1

0

一些随机的实验建议：

  • 稍后有障碍时应谨慎使用提前返回：工作组中的所有工作项都应执行障碍，如果某些工作项已经返回，您的内核将挂起。

  • 这种情况下的纹理缓存通常至少与本地内存一样快。尝试用简单的 read_imagef 调用替换 s0,...,s8 初始化。然后尝试不同的工作组维度。

你在什么硬件上运行？

于 2013-06-05T05:40:36.987 回答