optimization - 在 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 x 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

编辑 : 想想内存访问模式，原来的naive版本可能还不错；打电话时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 x n 像素的块，则边界将重叠 nx15(或 15xn)。您需要根据可用的本地内存大小 (LDS) 计算 n 的最大可能值。如果您使用的是 opencl 1.1 或更高版本，则 LDS 至少为 32kb。 opencl 1.0 promise 每个工作组 16kb。

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

使用 n=45 将使用 LDS 的 32768 个字节中的 32400 个，并让您每组使用 900 个工作项 (45-15)^2 = 900。例如 64x32 将使用所有 LDS，但组大小 = (64-15)*(32-15) = 833。

为内核使用 LDS 的步骤:

为图像的缓存块分配一维或二维本地数组。我使用 #define 常量，它很少需要更改。

从图像中读取 uint 值，并存储在本地。

调整每个工作项的 'pos' 以与本地内存相关

执行相同的 i,j 循环，但使用本地内存读取值。记住 i 和 j 循环在距离 n 15 处停止。

每个步骤如果不确定如何实现，可以在线搜索，或者您可以问我是否需要帮助。

您设备上的 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 for 循环，除了您省略 pos 和 pos0 值。仅使用 i 和 j 作为位置。以与最初相同的方式累积“差异”。

将解决方案写入'dest'

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

我希望这可以为您解决问题。