使用volatile内存实现OpenCL内核通信时NVIDIA平台挂起问题

我尝试实现两个OpenCL内核间的通信：一个工作内核（worker kernel）循环运行，另一个控制内核（control kernel）为其分配任务并通知结束。采用volatile设备缓冲区通信的方案在Intel OpenCL 2.1平台可正常运行，但在搭载Quadro P400的NVIDIA OpenCL 3.0 CUDA平台上程序会挂起，工作内核陷入无限循环。以下是最小复现示例（MWE），通过宏PLAT_IDX切换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;
}

问题根源 #

NVIDIA CUDA架构的OpenCL实现对全局内存的缓存机制是问题核心：

• 工作内核中的volatile __global uint *buf虽然标记了volatile，但NVIDIA的SM（流式多处理器）可能会将buf[0]的值缓存到寄存器或L1缓存中，并不会每次循环都去全局内存重新读取。
• 即使控制内核修改了buf[0]，工作内核无法感知到这个更新，从而一直停留在循环中。
• Intel平台的缓存一致性模型更宽松，volatile可以保证每次读取都从全局内存获取，因此能正常工作。

修复方案 #

要强制NVIDIA平台的内核每次都从全局内存读取最新值，有两种可靠方式：

方式1：添加全局内存栅栏 #

在工作内核的循环中插入内存栅栏，确保读取操作能获取到全局内存的最新值：

__kernel void loop(volatile __global uint *buf) {
    while (buf[0] != 123) {
        buf[1]++;
        // 确保全局内存的读写操作可见性
        mem_fence(CLK_GLOBAL_MEM_FENCE);
    }
}

方式2：使用原子读取操作 #

原子操作会绕过缓存直接访问全局内存，保证读取到最新值：

__kernel void loop(volatile __global uint *buf) {
    while (atomic_load(&buf[0]) != 123) {
        buf[1]++;
    }
}

更优实现方案 #

除了volatile缓冲区，还有几种更可靠的内核间通信方式：

• 原子操作+内存栅栏：兼容性最好，几乎所有OpenCL设备都支持，适合简单的控制信号传递。
• OpenCL 2.0信号量：如果设备支持OpenCL 2.0，可使用clCreateSemaphore、clEnqueueSignalSemaphore和clEnqueueWaitSemaphore实现内核间同步，比volatile缓冲区更可靠。
• 管道（Pipes）：虽支持设备较少，但如果目标平台支持，管道是专门为流式通信设计的，效率更高，适合任务分发场景。
• 主机作为中间层：若不需要严格的内核间并发，可让主机作为中介：工作内核完成一轮任务后返回主机，主机再分配新任务或通知结束，兼容性最好但延迟较高。

内容的提问来源于stack exchange，提问作者Changed My Name Again