Skip to content

Commit

Permalink
add an option for command buffer emulation enhanced error checking
Browse files Browse the repository at this point in the history 
Adds an option for enhanced error checking, disabled by default.
Enhanced error checking creates special test queues when a command
buffer is created, and enqueues a barrier blocked by a user event
into the test queue.  Then, before a command is recorded into a
command buffer, it is also enqueued into the test queue, to identify
command errors.  When the command buffer is finalized the user event
is set to an error state, causing all of the commands in the test
queue to be terminated.
  • Loading branch information
@bashbaug
bashbaug committed Jan 7, 2024
1 parent 7283fcb commit cb0659d
Show file tree
Hide file tree
Showing 3 changed files with 298 additions and 3 deletions.
287 changes: 286 additions & 1 deletion layers/10_cmdbufemu/emulate.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1224,9 +1224,13 @@ typedef struct _cl_command_buffer_khr
properties,
properties + numProperties );

cmdbuf->TestQueues.reserve(num_queues);
cmdbuf->BlockingEvents.reserve(num_queues);

for( auto queue : cmdbuf->Queues )
{
g_pNextDispatch->clRetainCommandQueue(queue);
cmdbuf->setupTestQueue(queue);
}
}

Expand All @@ -1239,6 +1243,19 @@ typedef struct _cl_command_buffer_khr
{
g_pNextDispatch->clReleaseCommandQueue(queue);
}

for( auto event : BlockingEvents )
{
g_pNextDispatch->clSetUserEventStatus(
event,
-1 );
g_pNextDispatch->clReleaseEvent(event);
}

for( auto queue : TestQueues )
{
g_pNextDispatch->clReleaseCommandQueue(queue);
}
}

static bool isValid( cl_command_buffer_khr cmdbuf )
Expand All @@ -1264,7 +1281,20 @@ typedef struct _cl_command_buffer_khr

cl_command_queue getQueue() const
{
return Queues[0];
if( Queues.size() > 0 )
{
return Queues[0];
}
return nullptr;
}

cl_command_queue getTestQueue() const
{
if( TestQueues.size() > 0 )
{
return TestQueues[0];
}
return nullptr;
}

cl_int getInfo(
Expand Down Expand Up @@ -1460,6 +1490,23 @@ typedef struct _cl_command_buffer_khr
return CL_INVALID_OPERATION;
}

for( auto event : BlockingEvents )
{
g_pNextDispatch->clSetUserEventStatus(
event,
-1 );
g_pNextDispatch->clReleaseEvent(event);
}

BlockingEvents.clear();

for( auto queue : TestQueues )
{
g_pNextDispatch->clReleaseCommandQueue(queue);
}

TestQueues.clear();

State = CL_COMMAND_BUFFER_STATE_EXECUTABLE_KHR;
return CL_SUCCESS;
}
Expand Down Expand Up @@ -1569,9 +1616,84 @@ typedef struct _cl_command_buffer_khr
cl_command_buffer_flags_khr Flags;
std::atomic<uint32_t> RefCount;

std::vector<cl_command_queue> TestQueues;
std::vector<cl_event> BlockingEvents;

std::vector<std::unique_ptr<Command>> Commands;
std::atomic<uint32_t> NextSyncPoint;

void setupTestQueue(cl_command_queue src)
{
if( g_cEnhancedErrorChecking )
{
cl_command_queue testQueue = nullptr;

cl_context context = nullptr;
g_pNextDispatch->clGetCommandQueueInfo(
src,
CL_QUEUE_CONTEXT,
sizeof(context),
&context,
nullptr );

cl_device_id device = nullptr;
g_pNextDispatch->clGetCommandQueueInfo(
src,
CL_QUEUE_DEVICE,
sizeof(device),
&device,
nullptr );

size_t propsSize = 0;
g_pNextDispatch->clGetCommandQueueInfo(
src,
CL_QUEUE_PROPERTIES_ARRAY,
0,
nullptr,
&propsSize );
if (propsSize != 0) {
size_t numProps = propsSize / sizeof(cl_queue_properties);
std::vector<cl_queue_properties> props(numProps);
g_pNextDispatch->clGetCommandQueueInfo(
src,
CL_QUEUE_PROPERTIES_ARRAY,
propsSize,
props.data(),
nullptr );
testQueue = g_pNextDispatch->clCreateCommandQueueWithProperties(
context,
device,
props.data(),
nullptr );
} else {
cl_command_queue_properties props = 0;
g_pNextDispatch->clGetCommandQueueInfo(
src,
CL_QUEUE_PROPERTIES,
sizeof(props),
&props,
nullptr );
testQueue = g_pNextDispatch->clCreateCommandQueue(
context,
device,
props,
nullptr );
}

cl_event blockingEvent = g_pNextDispatch->clCreateUserEvent(
context,
nullptr );
g_pNextDispatch->clEnqueueBarrierWithWaitList(
testQueue,
1,
&blockingEvent,
nullptr );

TestQueues.push_back(testQueue);
BlockingEvents.push_back(blockingEvent);
}
}

_cl_command_buffer_khr(cl_command_buffer_flags_khr flags) :
Magic(cMagic),
State(CL_COMMAND_BUFFER_STATE_RECORDING_KHR),
Expand Down Expand Up @@ -1784,6 +1906,23 @@ cl_int CL_API_CALL clCommandCopyBufferKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueCopyBuffer(
testQueue,
src_buffer,
dst_buffer,
src_offset,
dst_offset,
size,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}


cmdbuf->addCommand(
CopyBuffer::create(
Expand Down Expand Up @@ -1833,6 +1972,26 @@ cl_int CL_API_CALL clCommandCopyBufferRectKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueCopyBufferRect(
testQueue,
src_buffer,
dst_buffer,
src_origin,
dst_origin,
region,
src_row_pitch,
src_slice_pitch,
dst_row_pitch,
dst_slice_pitch,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
CopyBufferRect::create(
Expand Down Expand Up @@ -1882,6 +2041,22 @@ cl_int CL_API_CALL clCommandCopyBufferToImageKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueCopyBufferToImage(
testQueue,
src_buffer,
dst_image,
src_offset,
dst_origin,
region,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
CopyBufferToImage::create(
Expand Down Expand Up @@ -1927,6 +2102,22 @@ cl_int CL_API_CALL clCommandCopyImageKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueCopyImage(
testQueue,
src_image,
dst_image,
src_origin,
dst_origin,
region,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
CopyImage::create(
Expand Down Expand Up @@ -1972,6 +2163,23 @@ cl_int CL_API_CALL clCommandCopyImageToBufferKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueCopyImageToBuffer(
testQueue,
src_image,
dst_buffer,
src_origin,
region,
dst_offset,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}


cmdbuf->addCommand(
CopyImageToBuffer::create(
Expand Down Expand Up @@ -2017,6 +2225,22 @@ cl_int CL_API_CALL clCommandFillBufferKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueFillBuffer(
testQueue,
buffer,
pattern,
pattern_size,
offset,
size,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
FillBuffer::create(
Expand Down Expand Up @@ -2061,6 +2285,21 @@ cl_int CL_API_CALL clCommandFillImageKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueFillImage(
testQueue,
image,
fill_color,
origin,
region,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
FillImage::create(
Expand Down Expand Up @@ -2103,6 +2342,21 @@ cl_int CL_API_CALL clCommandSVMMemcpyKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueSVMMemcpy(
testQueue,
CL_FALSE,
dst_ptr,
src_ptr,
size,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
SVMMemcpy::create(
Expand Down Expand Up @@ -2145,6 +2399,21 @@ cl_int CL_API_CALL clCommandSVMMemFillKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueSVMMemFill(
testQueue,
dst_ptr,
pattern,
pattern_size,
size,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cmdbuf->addCommand(
SVMMemFill::create(
Expand Down Expand Up @@ -2190,6 +2459,22 @@ cl_int CL_API_CALL clCommandNDRangeKernelKHR_EMU(
{
return errorCode;
}
if( cl_command_queue testQueue = cmdbuf->getTestQueue() )
{
if( cl_int errorCode = g_pNextDispatch->clEnqueueNDRangeKernel(
testQueue,
kernel,
work_dim,
global_work_offset,
global_work_size,
local_work_size,
0,
nullptr,
nullptr ) )
{
return errorCode;
}
}

cl_int errorCode = CL_SUCCESS;
auto command = NDRangeKernel::create(
Expand Down
Loading

0 comments on commit cb0659d

Please sign in to comment.