使用易失性存储器的 OpenCL 内核通信

Question

我正在尝试让两个 OpenCL 内核相互通信。 A 工作内核运行一个循环，控制内核向其提供作业并告诉 当它完成时。我正在使用易失性设备缓冲区 沟通。当我使用 Intel OpenCL 2.1 平台时它可以工作， 但是当我使用 Nvidia OpenCL 3.0 CUDA (Quadro P400) 平台时 程序挂起。看来工作内核会永远循环。

MWE如下。

PLAT_IDX

定义可以是 0 或 1，并在 Intel 和 Nvidia 平台：

#include <assert.h>
#include <stdio.h>
#define CL_TARGET_OPENCL_VERSION 300
#include <CL/cl.h>
#define PLAT_IDX 0
#define CHK(x)      assert((x) == CL_SUCCESS)

const char *PROGRAM = (
    "__kernel void loop(volatile __global uint *buf) {"
    "   while (buf[0] != 123) { buf[1]++; }"
    "}"
    "__kernel void post(volatile __global uint *buf) {"
    "   buf[0] = 123;"
    "}"
);

int main(int argc, char *argv[]) {
    cl_uint n_platforms, n_devices;
    CHK(clGetPlatformIDs(0, NULL, &n_platforms));
    cl_platform_id plats[2];
    CHK(clGetPlatformIDs(n_platforms, plats, NULL));
    CHK(clGetDeviceIDs(plats[PLAT_IDX], CL_DEVICE_TYPE_ALL, 0, NULL, &n_devices));
    cl_device_id dev;
    CHK(clGetDeviceIDs(plats[PLAT_IDX], CL_DEVICE_TYPE_ALL, 1, &dev, NULL));
    assert(n_platforms > 0 && n_devices > 0);

    cl_int err;
    cl_context ctx = clCreateContext(NULL, 1, &dev, NULL, NULL, &err);
    CHK(err);
    cl_command_queue_properties props[] = {
        CL_QUEUE_PROPERTIES,
        CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, 0
    };
    cl_command_queue queue = clCreateCommandQueueWithProperties(
        ctx, dev, props, &err);
    CHK(err);

    cl_program prog = clCreateProgramWithSource(ctx, 1, &PROGRAM, NULL, &err);
    CHK(err);
    CHK(clBuildProgram(prog, 1, &dev, NULL, NULL, NULL));
    cl_kernel loop = clCreateKernel(prog, "loop", &err);
    CHK(err);
    cl_kernel post = clCreateKernel(prog, "post", &err);
    CHK(err);
    cl_mem mem = clCreateBuffer(
        ctx, CL_MEM_READ_WRITE, 2 * sizeof(cl_uint), NULL, &err);
    CHK(err);
    cl_event evs[2];
    CHK(clSetKernelArg(loop, 0, sizeof(cl_mem), &mem));
    CHK(clSetKernelArg(post, 0, sizeof(cl_mem), &mem));
    CHK(clEnqueueNDRangeKernel(
            queue, loop, 1, NULL, (size_t[]){1},
            NULL, 0, NULL, &evs[0]));
    CHK(clEnqueueNDRangeKernel(
            queue, post, 1, NULL, (size_t[]){1},
            NULL, 0, NULL, &evs[1]));
    printf("Waiting for kernels\n");
    CHK(clWaitForEvents(2, evs));
    return 1;
}

也许有更好的方法来实现这一点？ OpenCL 2.0 有管道， 但很少有设备支持它们。

Answer 1

您不能让两个 OpenCL 内核进行通信。易失性关键字也无法实现这一点。内核被放置在一个队列中并一个接一个地执行。仅在单独的队列中，内核可能会同时执行，但不能保证。

尝试在不需要内核通信的情况下解决它。让一个内核完成并将结果存储在全局内存中，然后另一个内核处理第一个内核的结果。

---

顺便说一句，您可以在一个内核内跨线程进行通信。只要线程在同一个工作组中，就可以使用本地内存。如果您需要跨工作组进行通信，您仍然可以使用原子来实现。