如何利用 OpenCL 中的 CL_MEM_ALLOC_HOST_PTR 标志来利用统一内存？

Question

如何在不正确复制数据的情况下在 OpenCL 设备上使用统一内存？

OpenCL 定义 CL_DEVICE_HOST_UNIFIED_MEMORY 为：

如果设备和主机具有统一的内存子系统，则为 CL_TRUE，否则为 CL_FALSE。

ARM Mali docs 建议使用统一内存，创建缓冲区时应使用标志

CL_MEM_ALLOC_HOST_PTR

。该标志描述为：

这是对驱动程序的提示，表明缓冲区在主机端被访问。要在应用程序处理器端使用缓冲区，必须映射该缓冲区并将数据写入其中。这是唯一不涉及复制数据的方法。如果必须填写由 GPU 处理的图像，这是避免复制的最佳方法。

但是，没有如何使用通过标志

CL_MEM_ALLOC_HOST_PTR

创建的缓冲区的示例。我写的例子似乎做得不太好。

考虑以下代码片段来使用此类缓冲区：

    // Create Buffers
    constexpr size_t n_bytes = sizeof(int) * SZ_ARR;
    cl::Buffer buffer_A(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
                        n_bytes);
    cl_int error{0};
    int* A = static_cast<int*>(
        queue.enqueueMapBuffer(buffer_A, CL_FALSE, CL_MAP_WRITE, 0, n_bytes,
                               nullptr, nullptr, &error));
    gpuErrchk(error);
    for (size_t i = 0; i < SZ_ARR; ++i) {
        A[i] = i;
    }
    gpuErrchk(queue.enqueueUnmapMemObject(buffer_A, A));
    cl::Kernel add(program, "add_and_print");
    add.setArg(0, buffer_A);
    gpuErrchk(queue.enqueueNDRangeKernel(add, cl::NullRange,
                                         cl::NDRange(SZ_ARR), cl::NullRange));
    queue.finish();

我正在创建缓冲区

buffer_A

以包含十个整数。我希望在使用

enqueueMapBuffer

映射它们然后使用

enqueueUnmapMemObject

取消映射之后写入这十个整数。内核

add_and_print

将

1

添加到每个数组元素并打印结果值。具体来说，内核是：

    const std::string kernel_code =
        "   void kernel add_and_print(global int* A) {"
        "       int i = get_global_id(0);"
        "       A[i] = A[i] + 1;"
        "       printf(\"%d\", A[i]);"
        "   }";

但是，程序会为每个数组元素打印

1

。

如何正确使用统一缓冲区？

作为参考，重现代码的完整程序如下：

#include <CL/opencl.hpp>
#include <iostream>

#define gpuErrchk(ans) \
    { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cl_int code, const char* file, int line,
                      bool abort = true) {
    if (code != CL_SUCCESS) {
        fprintf(stderr, "GPUassert, error code: %d %s %d\n", code, file, line);
        if (abort) exit(code);
    }
}

constexpr size_t SZ_ARR = 10;

cl::Device getDevice() {
    std::vector<cl::Platform> all_platforms;
    gpuErrchk(cl::Platform::get(&all_platforms));
    cl::Platform default_platform = all_platforms[0];
    std::vector<cl::Device> all_devices;
    gpuErrchk(default_platform.getDevices(CL_DEVICE_TYPE_GPU, &all_devices));
    cl::Device default_device = all_devices[0];
    std::cout << "Using device: " << default_device.getInfo<CL_DEVICE_NAME>()
              << "\n";
    return default_device;
}

cl::Program buildProgram(cl::Context& context, cl::Device& device) {
    const std::string kernel_code =
        "   void kernel add_and_print(global int* A) {"
        "       int i = get_global_id(0);"
        "       A[i] = A[i] + 1;"
        "       printf(\"%d\", A[i]);"
        "   }";
    cl::Program::Sources sources{{kernel_code.c_str(), kernel_code.length()}};

    cl::Program program(context, sources);
    if (program.build({device}) != CL_SUCCESS) {
        std::cout << "Error building: "
                  << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device)
                  << std::endl;
        exit(1);
    }
    return program;
}

int main() {
    // Prologue: Get device, context, and build program
    cl::Device default_device = getDevice();
    cl::Context context({default_device});
    cl::Program program = buildProgram(context, default_device);
    cl::CommandQueue queue(context, default_device);

    // Create Buffers
    constexpr size_t n_bytes = sizeof(int) * SZ_ARR;
    cl::Buffer buffer_A(context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
                        n_bytes);
    cl_int error{0};
    int* A = static_cast<int*>(
        queue.enqueueMapBuffer(buffer_A, CL_FALSE, CL_MAP_WRITE, 0, n_bytes,
                               nullptr, nullptr, &error));
    gpuErrchk(error);
    for (size_t i = 0; i < SZ_ARR; ++i) {
        A[i] = i;
    }
    gpuErrchk(queue.enqueueUnmapMemObject(buffer_A, A));
    cl::Kernel add(program, "add_and_print");
    add.setArg(0, buffer_A);
    gpuErrchk(queue.enqueueNDRangeKernel(add, cl::NullRange,
                                         cl::NDRange(SZ_ARR), cl::NullRange));
    queue.finish();
}

Answer 1

这确实是OpenCL统一内存的正确使用方式。程序中只有一个小错误，它没有等待

clEnqueueMapBuffer

完成。本质上第二个参数应该更改为

CL_TRUE

就像

int* A = static_cast<int*>(
    queue.enqueueMapBuffer(buffer_A, CL_TRUE, CL_MAP_WRITE, 0, n_bytes,
                           nullptr, nullptr, &error));

这应该会从 OpenCL 内核产生以下预期输出

$ clang++-15 -Iexternal/OpenCL-CLHPP/include -Iexternal/OpenCL-Headers/ -DCL_HPP_TARGET_OPENCL_VERSION=300 -DCL_TARGET_OPENCL_VERSION=300 -lOpenCL opencl_unified_memory.cpp -o opencl_um 
$ ./opencl_um 
Using device: NVIDIA T1200 Laptop GPU
12345678910

在具有许多设备内核的典型 OpenCL 程序中，应用程序代码可能会将多个缓冲区排入队列，而无需等待 OpenCL 运行时完成。然后，应用程序线程可以继续执行一些有意义的工作，然后通过检查

event

参数返回到映射缓冲区。请注意，通过阻止映射调用或通过事件引入的同步会降低 GPU 吞吐量，理想情况下应该在 OpenCL 设备上的计算周期结束时完成。