OpenCL 从内核发送数据到主机

问题描述

我试图制作一个内核，它偶尔会更新主机关于其进度的信息，因为该内核执行一个耗时的计算任务（需要几小时甚至几天才能完成）。然而，除了使用 printf 之外，我还没有找到一种从内核到主机的通信方式，这让我相信这种通信在理论上是可行的。

到目前为止，我尝试使用一个缓冲区，主机会定期检查是否有新的消息，但对于我的Nvidia GPU来说，在内核仍在执行时访问GPU内存的任何方式都是被锁定的。我尝试过使用共享内存，以及映射一个缓冲区，但都没有成功。

我也知道在OpenCL中存在事件，但我不确定如何在主机上创建一个用户事件并将其传递给内核。

我陷入了一个死胡同，任何帮助都将受到赞赏。我对OpenCL没有太多经验，所以我对大多数OpenCL的结构和概念不太了解。

这是我尝试过的代码，使用缓冲区：

#include <cstdio>
#include <CL/opencl.hpp>
#include <chrono>
#include <thread>

#define WORKER_COUNT 1024

struct packet_t {
    int completed_count;
    bool pending;
};

bool running = true;
cl_int err;
const char *kernel_code = "\
struct packet_t {\
    int completed_count;\
    bool pending;\
};\
__kernel void run(volatile __global struct packet_t *packet) {\
    size_t tid = get_global_id(0);\
    packet[tid].completed_count = 10;\
    packet[tid].pending = true;\
}";

void handle() {
    if (err != CL_SUCCESS) {
        fprintf(stderr, "Error %d!\n", err);
        exit(err);
    }
}

int main() {
    cl::Platform platform = cl::Platform::get(&err); handle(); // handle will check for errors
    cl_context_properties cprops[] = { CL_CONTEXT_PLATFORM, (cl_context_properties)platform(), 0 };
    cl::Context context { CL_DEVICE_TYPE_ALL, cprops, NULL, NULL, &err }; handle();
    cl::CommandQueue queue(context, context.getInfo<CL_CONTEXT_DEVICES>()[0], 0, &err); handle();
    cl::Event event;

    cl::Buffer packetsBuff { context, CL_MEM_ALLOC_HOST_PTR, sizeof(packet_t) * WORKER_COUNT, NULL, &err }; handle();
    volatile packet_t *packets = (packet_t*)queue.enqueueMapBuffer(packetsBuff, true, CL_MAP_READ, 0, sizeof(packet_t) * WORKER_COUNT, NULL, NULL, &err);

    cl::Program program { context, kernel_code };
    err = program.build(); handle();

    cl::Kernel kernel { program, "run", &err }; handle();
    err = kernel.setArg(0, packetsBuff); handle();
    err = queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(WORKER_COUNT), cl::NDRange(128), NULL, &event);
    handle();

    event.setCallback(CL_COMPLETE, [](cl_event e, cl_int i, void *v) { running = false; });

    while (running) {
        std::this_thread::sleep_for(std::chrono::milliseconds(100));
        for (size_t i = 0; i < WORKER_COUNT; i++) {
            if (packets[i].pending) {
                printf("Worker %d has done %d stuff.\n", packets[i].completed_count);
            }
        }
    }
}

解决方案

除非你有充分的理由以更复杂的方式进行操作，否则我强烈建议您只让内核运行一个很短的时间，比如100毫秒，然后将它所需要的任何内容保存到全局内存中。与局部内存或私有内存不同，全局内存在内核调用之间始终保持不变。我的意思是让内核按照以下方式工作。

A. Load state from global buffer
B. Do some work
C. Save state to global buffer

然后您的主机可以通过循环读取内核调用之间的缓冲区。