如何将设备端命令队列与主机端队列同步？ clFinish() 和 markerWithWaitList 给出无效队列错误

Question

我正在使用 OpenCL 2.0 动态并行功能，并让每个工作项将另一个内核与单个工作项排入队列。当子内核的工作完成时间较长时，父内核在子内核之前完成，并且不会保留内存一致性并返回损坏的数据（随机更新的数据项）。

由于 clFinish() 和 clEnqueueMarkerWithWaitList() 仅用于主机队列，因此我不能将它们用于此默认设备上的无序队列。

如何使子内核在某个同步点之前或至少在缓冲区读取命令之前完成，以实现内存一致性？

代码如下：

__kernel void test( __global float *xyz,__global float *xyzn,__global float *xyzo,__global float * arguments)
{
    int threadId=get_global_id(0);
    float dx=xyz[threadId*3]-arguments[2];float dy=xyz[threadId*3+1]-arguments[3];float t=arguments[1];
    float ctr=arguments[0];float wave=0.02f*ctr*sin(40.0f*t+100.0f*sqrt(dx*dx+dy*dy));
    xyzo[threadId*3]=xyz[threadId*3]+xyzn[threadId*3]*wave; // wave equation for all surface vertices
    xyzo[threadId*3+1]=xyz[threadId*3+1]+xyzn[threadId*3+1]*wave; // wave equation for all surface vertices
    xyzo[threadId*3+2]=xyz[threadId*3+2]+xyzn[threadId*3+2]*wave; // wave equation for all surface vertices
}

__kernel void waveEquation( __global float *xyz,__global float *xyzn,__global float *xyzo,__global float * arguments)
{
    int threadId=get_global_id(0);
    if(threadId<arguments[4])
    {
            queue_t q = get_default_queue();
            ndrange_t ndrange = ndrange_1D(threadId,1,1);
            void (^my_block_A)(void) = ^{test(xyz,xyzn,xyzo,arguments);};
            enqueue_kernel(q, CLK_ENQUEUE_FLAGS_NO_WAIT,ndrange,my_block_A);

    }

}

当父内核只有 1-2 个工作项时，它可以正常工作，但通常有 256*224 个工作项用于父内核，并且子内核在从主机访问数据之前无法完成（在 clFinish() 之后）

这里是默认队列的构造（不同于父内核的队列）

commandQueue = cl::CommandQueue(context, device,
   CL_QUEUE_ON_DEVICE|
   CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | 
   CL_QUEUE_ON_DEVICE_DEFAULT, &err);

编辑：这种创建队列的方式也不会使其可同步：

cl_uint qs=device.getInfo<CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE>();
cl_queue_properties qprop[] = { CL_QUEUE_SIZE, qs, CL_QUEUE_PROPERTIES, 
     (cl_command_queue_properties)(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE |
                                   CL_QUEUE_ON_DEVICE | 
                                   CL_QUEUE_ON_DEVICE_DEFAULT | 
                                   CL_QUEUE_PROFILING_ENABLE), 0 };
device_queue = clCreateCommandQueueWithProperties(context.get(),
                                   device.get(), qprop, &err);

设备=RX550，驱动程序=17.6.2，64 位构建。

---

User Parallel Highway 的解决方案也不起作用：

if(threadId<arguments[4])
{
        clk_event_t markerEvent;
        clk_event_t events[1];
        queue_t q = get_default_queue();
        ndrange_t ndrange = ndrange_1D(threadId,1,1);
        void (^my_block_A)(void) = ^{test(xyz,xyzn,xyzo,arguments);};
        enqueue_kernel(q, CLK_ENQUEUE_FLAGS_NO_WAIT,ndrange,0,NULL,&events[0],my_block_A);
        enqueue_marker(q, 1, events, &markerEvent);

        release_event(events[0]);
        release_event(markerEvent);

}

---

这不起作用：

queue_t q = get_default_queue();
ndrange_t ndrange = ndrange_1D(threadId,1,1);
void (^my_block_A)(void) = ^{test(xyz,xyzn,xyzo,arguments);};
int ctr=0;
while((enqueue_kernel(q, CLK_ENQUEUE_FLAGS_NO_WAIT,ndrange,my_block_A)&
        (   CLK_DEVICE_QUEUE_FULL|
            CLK_EVENT_ALLOCATION_FAILURE|
            CLK_OUT_OF_RESOURCES |
            CLK_INVALID_NDRANGE |
            CLK_INVALID_QUEUE |
            CLK_INVALID_EVENT_WAIT_LIST |
            CLK_INVALID_ARG_SIZE
        ))>0 )
{
}

这不起作用但完成，因此没有无限循环。

Answer 1

您应该考虑使用 enqueue_marker：

https://www.khronos.org/registry/OpenCL/specs/opencl-2.0-openclc.pdf#page=172

规范中还有一个示例，其中多个内核被排队，使用 enqueue_marker 命令，您可以等待子内核完成，然后继续处理父内核。示例代码在这里：

https://www.khronos.org/registry/OpenCL/specs/opencl-2.0-openclc.pdf#page=175

编辑：经过多次实验，结果如下： 随着父内核启动的子内核数量的增加，程序会失败。这可能是由于 huseyin tugrul buyukisik 建议的 queue_size 引起的。虽然执行没有返回错误代码，但结果是不正确的。 OpenCL 规范中没有提及此类问题。