使用填充的本地缓冲区 OpenCL 时像素值错误答案

【问题标题】：Wrong pixel values when using padded local buffer OpenCL使用填充的本地缓冲区 OpenCL 时像素值错误
【发布时间】：2023-03-16 06:39:01
【问题描述】：

当我使用本地缓冲区复制 OpenCL 内核中的数据时，我遇到了意外的结果。此处提供的代码非常简单（而且无用，因为我不需要使用本地缓冲区进行此类操作），但这是类似卷积过程的第一步。

这是我的代码：

std::string implementCopyFromLocalKernel()
{
    return BOOST_COMPUTE_STRINGIZE_SOURCE(
    __kernel void copyFromLocal_knl(__global const float* in,
                                    const ulong sizeX, const ulong sizeY,
                                    const int filterRadiusX, const int filterRadiusY,
                                    __local float* localImage,
                                    const ulong localSizeX, const ulong localSizeY,
                                    __global float* out)
    {
        // Store each work-item’s unique row and column
        const int x = get_global_id(0);
        const int y = get_global_id(1);

        // Group size
        int groupSizeX = get_local_size(0);
        int groupSizeY = get_local_size(1);

        // Determine the size of the work group output region
        int groupIdX = get_group_id(0);
        int groupIdY = get_group_id(1);

        // Determine the local ID of each work item
        int localX = get_local_id(0);
        int localY = get_local_id(1);

        // Padding
        int paddingX = filterRadiusX;
        int paddingY = filterRadiusY;

        // Cache the data to local memory
        // Copy the data for the current coordinates
        localImage[localX + localY*localSizeX] = in[x + y * sizeX];

        barrier(CLK_LOCAL_MEM_FENCE);

        out[x + y * sizeX] = localImage[localX + localY*localSizeX];

        return;
    }
    );
}

void copyLocalBuffer(const boost::compute::context& context,  boost::compute::command_queue& queue, const boost::compute::buffer& bufInn boost::compute::buffer& bufOut, const size_t sizeX, const size_t sizeY)
{
    const size_t nbPx = sizeX * sizeY;
    const size_t maxSize = (sizeX > sizeY ? sizeX : sizeY);

    // Prepare to launch the kernel
    std::string kernel_src = implementCopyFromLocalKernel();
    boost::compute::program program;
    try {
        program = boost::compute::program::create_with_source(kernel_src, pGpuDescription->getContext(deviceIdx));
        program.build();
    }
    catch (const boost::compute::opencl_error& e) {
        std::cout << "Error bulding program from source : " << std::endl << e.what() << std::endl
            << program.build_log() << std::endl;
        return;
    }

    boost::compute::kernel kernel;
    try {
        kernel = program.create_kernel("copyFromLocal_knl");
    }
    catch (const boost::compute::opencl_error& e) {
        std::cout << "Error creating kernel : " << std::endl << e.what() << std::endl;
        return;
    }

    try {
        int localSizeX = 16;
        int localSizeY = 16;
        int paddingPixelsX = 2;// 0; // <- Changing to 0 works
        int paddingPixelsY = paddingPixelsX;

        int localWidth = localSizeX + 2 * paddingPixelsX;
        int localHeight = localSizeY + 2 * paddingPixelsY;

        boost::compute::buffer localImage(context, localWidth*localHeight * sizeof(float));

        kernel.set_arg(0, bufIn);
        kernel.set_arg(1, sizeX);
        kernel.set_arg(2, sizeY);
        kernel.set_arg(3, paddingPixelsX);
        kernel.set_arg(4, paddingPixelsY);
        kernel.set_arg(5, localImage);
        kernel.set_arg(6, localWidth);
        kernel.set_arg(7, localHeight);
        kernel.set_arg(8, bufOut);
    }
    catch (const boost::compute::opencl_error& e) {
        std::cout << "Error setting kernel arguments: " << std::endl << e.what() << std::endl;
        return;
    }

    try {

        size_t origin[2] = { 0, 0 };
        size_t region[2] = { 256, 256 };// { sizeX, sizeY };
        size_t localSize[2] = { 16, 16 };
        queue.enqueue_nd_range_kernel(kernel, 2, origin, region, localSize);
    }
    catch (const boost::compute::opencl_error& e) {
        std::cout << "Error executing kernel : " << std::endl << e.what() << std::endl;
        return;
    }
}

我将代码简化为简单地复制与本地图像相关的本地坐标中的每个工作项对应的像素。因此，本地图像缓冲区的每一行都必须有2*paddingPixelsX 和2*paddingPixelsY 未使用行的未使用数据。

如果我不添加填充数据（paddingPixelsX 和paddingPixelsY = 0）它可以工作，但似乎某些工作项不会从输入缓冲区读取数据或将数据写入输出缓冲区（或本地缓冲区？）在正确的位置。而且，当我多次运行我的程序时，我从来没有得到相同的结果。

这是我得到的结果示例（右），山魈图像作为输入（左）：

我确保线程与barrier(CLK_LOCAL_MEM_FENCE); 同步，并且每个工作项读取和写入特定数据，如果我的代码有错误，我不明白为什么没有填充不会出错。

有人有想法吗？

谢谢，

【问题讨论】：

我不确定在许多工作组上下文中传递本地内存缓冲区作为内核参数是否是正确的方法，因为似乎只有一个本地缓冲区实例被传递。您可以尝试在内核中实例化本地缓冲区吗？例如：__local float localImage[16*16];.
感谢您的帮助！我将缓冲区作为参数传递，因为我事先不知道填充，所以我不能静态分配缓冲区。不过，您指出了问题：所有工作组都使用相同的缓冲区。我需要为每个工作组设置一个专用缓冲区。

标签： c++ image-processing boost parallel-processing opencl

【解决方案1】：

正如已经确认的那样，问题在于传递给内核的动态分配的本地缓冲区仅为一个工作组创建。

其中一种解决方案是在内核内部静态创建本地缓冲区，例如： __local float localImage[16*16];

如果缓冲区的大小不能被硬编码，那么它可以通过预处理器设置： __local float localImage[SIZE_X*SIZE_Y]; 然后这些参数在内核构建期间传递。根据我的记忆，使用内核参数定义静态本地缓冲区的大小可能不适用于每个 GPU（编译将失败）。

我不熟悉 boost 计算，但我认为应该可以通过将参数传递给 implementCopyFromLocalKernel() 来实现类似的效果，然后它们将在字符串化期间转换为值。

【讨论】：

感谢您的建议。只要性能不受影响，我更喜欢限制构建选项。我也可以直接附加到内核源字符串中，但代码会变得更难阅读。使用clSetKernelArg()的定义对我来说更自然。

【解决方案2】：

感谢@doqtor，我了解到问题来自作为内核参数传递的缓冲区。因此，所有工作组都使用相同的缓冲区。

由于我不知道卷积操作所需的填充大小，我需要这个缓冲区作为参数。我修改了内核参数化，以便每个工作组使用不同的缓冲区：

kernel.set_arg(5, localWidth*localHeight*sizeof(float), NULL);

我在阅读clSetKernelArg的documentation时错过了重要的部分：

如果使用 __local 限定符声明参数，则 arg_value 条目必须为 NULL。

【讨论】：