在设备和主机之间来回复制全局设备上指针地址答案

【问题标题】：Copying global on-device pointer address back and forth between device and host在设备和主机之间来回复制全局设备上指针地址
【发布时间】：2016-03-30 00:22:12
【问题描述】：

我在 OpenCL 设备（GPU）上创建了一个缓冲区，并且我需要从主机知道全局设备上指针地址，以便我可以将该设备上地址放在另一个缓冲区中，以便内核可以从包含第一个缓冲区地址的缓冲区中读取，以便它可以访问该缓冲区的内容。

如果这令人困惑，这就是我正在尝试做的事情：我创建一个表示 2D 图像的通用浮点缓冲区，然后从主机创建一个待办事项列表，其中列出了我的内核需要绘制的所有内容，哪些行，哪些圆圈，哪些图像...因此，内核必须从该列表中知道在哪里可以找到该图像，但是对该图像的引用不能作为内核参数传递，因为该内核可能不会绘制任何图像或一千个不同的图像，这一切都取决于列表所说的内容，因此必须在用作我的内核的待办事项列表的缓冲区中引用它。

到目前为止我做这件事的尴尬方式：

为此，我尝试创建一个函数，该函数在创建图像缓冲区后调用内核，该缓冲区获取缓冲区并将全局设备上地址作为 ulong 返回另一个缓冲区，然后主机将该值存储在 64 -bit 整数，像这样：

uint64_t get_clmem_device_address(clctx_t *clctx, cl_mem buf)
{
    const char kernel_source[] =
"kernel void get_global_ptr_address(global void *ptr, global ulong *devaddr)        \n"
"{                                          \n"
"   *devaddr = (ulong) ptr;                             \n"
"}                                          \n";

    int32_t i;
    cl_int ret;
    static int init=1;
    static cl_program program;
    static cl_kernel kernel;
    size_t global_work_size[1];
    static cl_mem ret_buffer;
    uint64_t devaddr;

    if (init)
    {
        init=0;
        ret = build_cl_program(clctx, &program, kernel_source);
        ret = create_cl_kernel(clctx, program, &kernel, "get_global_ptr_address");
        ret_buffer = clCreateBuffer(clctx->context, CL_MEM_WRITE_ONLY, 1*sizeof(uint64_t), NULL, &ret);
    }
    if (kernel==NULL)
        return ;

    // Run the kernel
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &buf);
    ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &ret_buffer);

    global_work_size[0] = 1;
    ret = clEnqueueNDRangeKernel(clctx->command_queue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);         // enqueue the kernel
    ret = clEnqueueReadBuffer(clctx->command_queue, ret_buffer, CL_FALSE, 0, 1*sizeof(uint64_t), &devaddr, 0, NULL, NULL);      // copy the value
    ret = clFlush(clctx->command_queue);
    clFinish(clctx->command_queue);

    return devaddr;
}

显然这行得通（它确实返回一个数字，虽然很难知道它是否正确），但是我把这个devaddr（主机上的一个 64 位整数）放在内核使用的待办事项列表缓冲区中知道要做什么，然后如果需要（根据列表）内核调用下面的函数，le 这里是指向待办事项列表中相关条目的指针，64 位地址是第一个元素：

float4 blit_sprite(global uint *le, float4 pv)
{
    const int2 p = (int2) (get_global_id(0), get_global_id(1));
    ulong devaddr;
    global float4 *im;
    int2 im_dim;

    devaddr = ((global ulong *) le)[0];     // global address for the start of the image as a ulong
    im_dim.x = le[2];
    im_dim.y = le[3];

    im = (global float4 *) devaddr;     // ulong is turned into a proper global pointer

    if (p.x < im_dim.x)
        if (p.y < im_dim.y)
            pv += im[p.y * im_dim.x + p.x];     // this gives me a CL_OUT_OF_RESOURCES error, even when changing it to im[0]

    return pv;
}

~~但令人惊讶的是，这不起作用，它给了我一个 CL_OUT_OF_RESOURCES 错误，我认为这意味着我的 im 指针无效。~~ 实际上它有效，它在我使用了两种不同的上下文。但它仍然很笨拙。

有没有一种不那么奇怪的方式来做我想做的事？

【问题讨论】：

标签： opencl

【解决方案1】：

OpenCL 标准不保证内存对象不会在内核调用之间进行物理重新分配。因此，原始设备端地址仅在单个内核 NDRange 内有效。这就是为什么 OpenCL 内存对象在 Host 端被表示为透明结构指针的原因之一。

不过，您可以将偏移量保存到第一个内核中内存对象的第一个字节并将其传递给第二个内核。每次启动内核时，您都会在内核中获得实际的设备端地址，并通过保存的移位值增加它。那将是完全“合法的”。

【讨论】：