使用填充的本地缓冲区 OpenCL

如何解决使用填充的本地缓冲区 OpenCL

当我使用本地缓冲区复制 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];

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

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

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

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

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

我在阅读clSetKernelArg的{{3}}时错过了重要的部分：

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

使用填充的本地缓冲区 OpenCL

如何解决使用填充的本地缓冲区 OpenCL

解决方法

相关推荐