在OpenCL中对偏移邻域构建操作的更快方法

Question

如何在2D阵列的许多重叠但偏移的块上进行操作,以便在OpenCL中更有效地执行？

例如,我有以下OpenCL内核:

__kernel void test_kernel(
    read_only image2d_t src,
    write_only image2d_t dest,
    const int width,
    const int height
)
{
    const sampler_t sampler =  CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    int2 pos = (int2)(get_global_id(0), get_global_id(1));
    int2 pos0 = (int2)(pos.x - pos.x % 16, pos.y - pos.y % 16);

    uint4 diff = (uint4)(0, 0, 0, 0);

    for (int i=0; i<16; i++)
    {
        for (int j=0; j<16; j++)
        {
            diff += read_imageui(src, sampler, (int2)(pos0.x + i, pos0.y + j)) -
                read_imageui(src, sampler, (int2)(pos.x + i, pos.y + j));
        }
    }
    write_imageui(dest, pos, diff);
}

它可以产生正确的结果,但速度很慢...... NVS4200M只有~25 GFLOPS,1k×1k输入.(硬件规格为155 GFLOPS).我猜这与内存访问模式有关.每个工作项读取一个16x16数据块,该数据块与16x16区域中的所有邻居相同,并且大多数时间的另一个偏移数据块与其直接邻居的数据重叠.所有读取都是通过采样器进行的 主机程序是PyOpenCL(我认为实际上没有改变任何东西),工作组大小是16x16.

编辑:下面的每个建议的新版本的内核,将工作区复制到本地变量:

__kernel __attribute__((reqd_work_group_size(16, 16, 1)))
void test_kernel(
    read_only image2d_t src,
    write_only image2d_t dest,
    const int width,
    const int height
)
{
    const sampler_t sampler =  CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;
    int2 pos = (int2)(get_global_id(0), get_global_id(1));

    int dx = pos.x % 16;
    int dy = pos.y % 16;

    __local uint4 local_src[16*16];
    __local uint4 local_src2[32*32];

    local_src[(pos.y % 16) * 16 + (pos.x % 16)] = read_imageui(src, sampler, pos);
    local_src2[(pos.y % 16) * 32 + (pos.x % 16)] = read_imageui(src, sampler, pos);
    local_src2[(pos.y % 16) * 32 + (pos.x % 16) + 16] = read_imageui(src, sampler, (int2)(pos.x + 16, pos.y));
    local_src2[(pos.y % 16 + 16) * 32 + (pos.x % 16)] = read_imageui(src, sampler, (int2)(pos.x, pos.y + 16));
    local_src2[(pos.y % 16 + 16) * 32 + (pos.x % 16) + 16] = read_imageui(src, sampler, (int2)(pos.x + 16, pos.y + 16));
    barrier(CLK_LOCAL_MEM_FENCE);


    uint4 diff = (uint4)(0, 0, 0, 0);

    for (int i=0; i<16; i++)
    {
        for (int j=0; j<16; j++)
        {
            diff += local_src[ j*16 + i ] - local_src2[ (j+dy)*32 + i+dx ];
        }
    }
    write_imageui(dest, pos, diff);
}

结果:输出正确,运行时间慢56%.如果仅使用local_src(而不是local_src2),结果会快〜10%.

编辑:基于更强大的硬件基准测试,AMD Radeon HD 7850获得420GFLOPS,规格为1751GFLOPS.为了公平起见,规范是乘法加法,并且这里没有乘法因此预期为~875GFLOPS,但与理论性能相比,这仍然相当多.

编辑:为了方便运行测试的任何人想尝试这个,下面的PyOpenCL中的主机端程序:

import pyopencl as cl
import numpy
import numpy.random
from time import time

CL_SOURCE = ''' 
// kernel goes here
'''

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx, properties=cl.command_queue_properties.PROFILING_ENABLE)
prg = cl.Program(ctx, CL_SOURCE).build()

h, w = 1024, 1024
src = numpy.zeros((h, w, 4), dtype=numpy.uint8)
src[:,:,:] = numpy.random.rand(h, w, 4) * 255

mf = cl.mem_flags
src_buf = cl.image_from_array(ctx, src, 4)
fmt = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.UNSIGNED_INT8)
dest_buf = cl.Image(ctx, mf.WRITE_ONLY, fmt, shape=(w, h))

# warmup
for n in range(10):
    event = prg.test_kernel(queue, (w, h), (16,16), src_buf, dest_buf, numpy.int32(w), numpy.int32(h))
event.wait()

# benchmark
t1 = time()
for n in range(100):
    event = prg.test_kernel(queue, (w, h), (16,16), src_buf, dest_buf, numpy.int32(w), numpy.int32(h))
event.wait()
t2 = time()
print "Duration (host): ", (t2-t1)/100
print "Duration (event): ", (event.profile.end-event.profile.start)*1e-9

编辑:考虑内存访问模式,原始的天真版本可能相当不错; 当调用read_imageui(src, sampler, (int2)(pos0.x + i, pos0.y + j))工作组中的所有工作项时,读取相同的位置(所以这只是一个读取??),并且在调用时read_imageui(src, sampler, (int2)(pos.x + i, pos.y + j))它们正在读取顺序位置(因此读取可以完美地合并？).

Answer 1

这绝对是一个内存访问问题.相邻工作项的像素可以重叠15x16,更糟糕的是,每个工作项至少重叠225个.

我会使用本地内存并让工作组协同处理许多16x16块.我喜欢为每个工作组使用一个大的方块.矩形块有点复杂,但可以为您获得更好的内存利用率.

如果从源图像中读取n×n像素的块,则边界将重叠nx15(或15xn).您需要根据可用的本地内存大小(LDS)计算n的最大可能值.如果您使用的是opencl 1.1或更高版本,则LDS至少为32kb.opencl 1.0承诺每个工作组16kb.

n <= sqrt(32kb / sizeof(uint4))
n <= sqrt(32768 / 16)
n ~ 45

使用n = 45将使用327S字节的LDS中的32400,并允许每组使用900个工作项(45-15)^ 2 = 900.注意:这是矩形块有用的地方; 例如,64x32将使用所有LDS,但组大小=(64-15)*(32-15)= 833.

将LDS用于内核的步骤:

1. 为映像的缓存块分配1D或2D本地数组.我使用#define常量,很少需要改变.
2. 从图像中读取uint值,然后在本地存储.
3. 调整每个工作项的'pos'以与本地记忆相关
4. 执行相同的i,j循环,但使用本地内存读取值.记住,i和j循环停止15短于n.

如果您不确定如何实施,可以在线搜索每个步骤,或者您可以询问我是否需要帮忙.

设备上的LDS优于纹理读取速度的可能性很大.这是违反直觉的,但请记住,您一次只能读取少量数据,因此gpu可能无法有效地缓存像素.LDS的使用将保证像素可用,并且考虑到每个像素的读取次数,我希望这会产生巨大的差异.

请告诉我您观察到的结果.

更新:这是我试图更好地解释我的解决方案.我使用方格纸作为我的绘图,因为我对图像处理软件并不是那么好.

上面是在第一个代码片段中如何从src读取值的草图.最大的问题是pos0矩形--16x16 uint4值 - 正在为组中的每个工作项(其中256个)完整地读取.我的解决方案涉及读取大面积区域并共享所有256个工作组的数据.

如果将图像的31x31区域存储在本地存储器中,则所有256个工作项的数据都将可用.

脚步:

• 使用工作组维度:(16,16)
• 将src的值读入一个大的本地缓冲区,即:uint4 buff [31] [31]; 需要翻译缓冲区,使'pos0'处于buff [0] [0]
• 屏障(CLK_LOCAL_MEM_FENCE)等待内存复制操作
• 做同样的i,j你原来的循环,除了你省略pos和pos0值.只使用i和j作为位置.以与最初相同的方式累积'diff'.
• 把解决方案写成'dest'

这与我对您的问题的第一个回答相同,除了我使用n = 16.此值不会完全利用本地内存,但可能适用于大多数平台.256往往是一个共同的最大工作组大小.

我希望这能为你解决问题.