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

Question

当我使用本地缓冲区在 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); 同步，并且每个工作项都读取和写入特定数据，如果我的代码有问题，我不明白为什么没有填充不会出错。

有人有想法吗？

谢谢，

Answer 1

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

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

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

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

If the argument is declared with the __local qualifier, the arg_value entry must be NULL.

Answer 2

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

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

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

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

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

Wrong pixel values when using padded local buffer OpenCL

c++

image-processing

boost

parallel-processing

opencl