From cb0659d73df091b122170ca88e704c4d15925ea7 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Sat, 6 Jan 2024 17:11:39 -0800 Subject: [PATCH] add an option for command buffer emulation enhanced error checking 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. --- layers/10_cmdbufemu/emulate.cpp | 287 +++++++++++++++++++++++++++++++- layers/10_cmdbufemu/emulate.h | 6 +- layers/10_cmdbufemu/main.cpp | 8 + 3 files changed, 298 insertions(+), 3 deletions(-) diff --git a/layers/10_cmdbufemu/emulate.cpp b/layers/10_cmdbufemu/emulate.cpp index d45c526..efe2627 100644 --- a/layers/10_cmdbufemu/emulate.cpp +++ b/layers/10_cmdbufemu/emulate.cpp @@ -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); } } @@ -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 ) @@ -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( @@ -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; } @@ -1569,9 +1616,84 @@ typedef struct _cl_command_buffer_khr cl_command_buffer_flags_khr Flags; std::atomic RefCount; + std::vector TestQueues; + std::vector BlockingEvents; + std::vector> Commands; std::atomic 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 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), @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( @@ -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( diff --git a/layers/10_cmdbufemu/emulate.h b/layers/10_cmdbufemu/emulate.h index ea37874..bd672c9 100644 --- a/layers/10_cmdbufemu/emulate.h +++ b/layers/10_cmdbufemu/emulate.h @@ -9,6 +9,10 @@ #include +extern const bool g_cEnhancedErrorChecking; + +extern const struct _cl_icd_dispatch* g_pNextDispatch; + struct SLayerContext { typedef std::map CEventMap; @@ -17,8 +21,6 @@ struct SLayerContext SLayerContext& getLayerContext(void); -extern const struct _cl_icd_dispatch* g_pNextDispatch; - /////////////////////////////////////////////////////////////////////////////// // Emulated Functions diff --git a/layers/10_cmdbufemu/main.cpp b/layers/10_cmdbufemu/main.cpp index 07e9173..ee4cf05 100644 --- a/layers/10_cmdbufemu/main.cpp +++ b/layers/10_cmdbufemu/main.cpp @@ -27,6 +27,14 @@ #include "emulate.h" +// Enhanced error checking can be used to catch additional errors when +// commands are recorded into a command buffer, but relies on tricky +// use of user events that may not work properly with some implementations. +// Disabling enhanced error checkgin may enable command buffer emulation +// to function properly on more implementations. + +const bool g_cEnhancedErrorChecking = false; + const struct _cl_icd_dispatch* g_pNextDispatch = NULL; static cl_int CL_API_CALL