Как я могу синхронизировать очереди команд на стороне устройства с очередями на стороне хоста? 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())

Здесь строится очередь по умолчанию (отличная от очереди для parent-kernel)

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

В спецификации также есть пример, когда несколько ядер ставятся в очередь, и с помощью команды enqueue_marker вы можете дождаться завершения дочерних ядер, а затем продолжить работу с родительским ядром. Пример кода здесь:

https://www.khronos.org/registry/OpenCL/specs/opencl-2.0-openclc.pdf

Изменить: После нескольких экспериментов, выводы заключаются в следующем: По мере увеличения числа дочерних ядер, запускаемых родительским ядром, происходит сбой программы. Это, вероятно, вызвано queue_size, как предположил huseyin tugrul buyukisik. Хотя выполнение не возвращает код ошибки, результаты неверны. В спецификации OpenCL нет упоминания об этом типе проблемы.