如何将设备端命令队列与主机端队列同步? clFinish() 和 markerWithWaitList 给出无效队列错误
How can I synchronize device-side command queues with host-side queues? clFinish() and markerWithWaitList gives invalid queue error
我正在使用 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 位版本。
用户 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 )
{
}
这不起作用但会完成,因此没有无限循环。
你应该考虑使用 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 引起的。尽管执行没有 return 错误代码,但结果不正确。 OpenCL规范中没有提及此类问题。
我正在使用 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 位版本。
用户 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 )
{
}
这不起作用但会完成,因此没有无限循环。
你应该考虑使用 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 引起的。尽管执行没有 return 错误代码,但结果不正确。 OpenCL规范中没有提及此类问题。