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