我正在使用 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 )
{
}
这不起作用但完成，因此没有无限循环。