From cb0659d73df091b122170ca88e704c4d15925ea7 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Sat, 6 Jan 2024 17:11:39 -0800 Subject: [PATCH 1/4] 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 d45c5267..efe26276 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 ea37874f..bd672c9a 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 07e9173c..ee4cf05e 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 From d303ca57a9d05c65900f77b74766d1c23705f31f Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Sat, 6 Jan 2024 17:16:21 -0800 Subject: [PATCH 2/4] fix typo --- layers/10_cmdbufemu/main.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/layers/10_cmdbufemu/main.cpp b/layers/10_cmdbufemu/main.cpp index ee4cf05e..d30f0a49 100644 --- a/layers/10_cmdbufemu/main.cpp +++ b/layers/10_cmdbufemu/main.cpp @@ -30,7 +30,7 @@ // 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 +// Disabling enhanced error checking may enable command buffer emulation // to function properly on more implementations. const bool g_cEnhancedErrorChecking = false; From dea60984b07e2edf9c3ec247390a4817df5ed095 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Sun, 26 May 2024 10:29:42 -0700 Subject: [PATCH 3/4] add an environment variable to control enhanced error checking --- include/getenv_util.hpp | 98 +++++++++++++++++++++++++++++++++ layers/10_cmdbufemu/README.md | 10 +++- layers/10_cmdbufemu/emulate.cpp | 2 +- layers/10_cmdbufemu/emulate.h | 2 +- layers/10_cmdbufemu/main.cpp | 7 ++- 5 files changed, 113 insertions(+), 6 deletions(-) create mode 100644 include/getenv_util.hpp diff --git a/include/getenv_util.hpp b/include/getenv_util.hpp new file mode 100644 index 00000000..9e24f929 --- /dev/null +++ b/include/getenv_util.hpp @@ -0,0 +1,98 @@ +/* +// Copyright (c) 2022-2024 Ben Ashbaugh +// +// SPDX-License-Identifier: MIT +*/ +#pragma once + +#include +#include + +#include + +#if defined(_WIN32) + +#include + +#define GETENV( _name, _value ) _dupenv_s( &_value, NULL, _name ) +#define FREEENV( _value ) free( _value ) + +#else + +#define GETENV( _name, _value ) _value = getenv(_name) +#define FREEENV( _value ) (void)_value + +#endif + +static inline bool getControlFromEnvironment( + const char* name, + void* pValue, + size_t size ) +{ + char* envVal = NULL; + GETENV( name, envVal ); + + if( envVal != NULL ) + { + if( size == sizeof(unsigned int) ) + { + unsigned int* puVal = (unsigned int*)pValue; + *puVal = atoi(envVal); + } + else if( strlen(envVal) < size ) + { + char* pStr = (char*)pValue; + strcpy( pStr, envVal ); + } + + FREEENV( envVal ); + return true; + } + + return false; +} + +template +static bool getControl( + const char* name, + T& value ) +{ + unsigned int readValue = 0; + bool success = getControlFromEnvironment( name, &readValue, sizeof(readValue) ); + if( success ) + { + value = readValue; + } + + return success; +} + +template <> +bool getControl( + const char* name, + bool& value ) +{ + unsigned int readValue = 0; + bool success = getControlFromEnvironment( name, &readValue, sizeof(readValue) ); + if( success ) + { + value = ( readValue != 0 ); + } + + return success; +} + +template <> +bool getControl( + const char* name, + std::string& value ) +{ + char readValue[256] = ""; + bool success = getControlFromEnvironment( name, readValue, sizeof(readValue) ); + if( success ) + { + value = readValue; + } + + return success; +} diff --git a/layers/10_cmdbufemu/README.md b/layers/10_cmdbufemu/README.md index 169b47e0..418af38a 100644 --- a/layers/10_cmdbufemu/README.md +++ b/layers/10_cmdbufemu/README.md @@ -26,9 +26,17 @@ clGetExtensionFunctionAddressForPlatform clInitLayer ``` +## Optional Controls + +The following environment variables can modify the behavior of the command buffer emulation layer: + +| Environment Variable | Behavior | Example Format | +|----------------------|----------|-----------------| +| `CMDBUFEMU_EnhancedErrorChecking` | Enables additional error checking when commands are added to a command buffer using a command buffer "test queue". By default, the additional error checking is disabled. | `export CMDBUFEMU_EnhancedErrorChecking=1`

`set CMDBUFEMU_EnhancedErrorChecking=1` | + ## Known Limitations This section describes some of the limitations of the emulated `cl_khr_command_buffer` functionality: -* Many error conditions are not properly checked for and returned. +* Some error conditions are not properly checked for and returned. * Many functions are not thread safe. diff --git a/layers/10_cmdbufemu/emulate.cpp b/layers/10_cmdbufemu/emulate.cpp index 7e389d40..fe2f4454 100644 --- a/layers/10_cmdbufemu/emulate.cpp +++ b/layers/10_cmdbufemu/emulate.cpp @@ -1645,7 +1645,7 @@ typedef struct _cl_command_buffer_khr void setupTestQueue(cl_command_queue src) { - if( g_cEnhancedErrorChecking ) + if( g_EnhancedErrorChecking ) { cl_command_queue testQueue = nullptr; diff --git a/layers/10_cmdbufemu/emulate.h b/layers/10_cmdbufemu/emulate.h index bd672c9a..15922fda 100644 --- a/layers/10_cmdbufemu/emulate.h +++ b/layers/10_cmdbufemu/emulate.h @@ -9,7 +9,7 @@ #include -extern const bool g_cEnhancedErrorChecking; +extern bool g_EnhancedErrorChecking; extern const struct _cl_icd_dispatch* g_pNextDispatch; diff --git a/layers/10_cmdbufemu/main.cpp b/layers/10_cmdbufemu/main.cpp index 817a15f7..1194b962 100644 --- a/layers/10_cmdbufemu/main.cpp +++ b/layers/10_cmdbufemu/main.cpp @@ -23,6 +23,7 @@ #include #include +#include "getenv_util.hpp" #include "layer_util.hpp" #include "emulate.h" @@ -30,10 +31,8 @@ // 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 checking may enable command buffer emulation -// to function properly on more implementations. -const bool g_cEnhancedErrorChecking = true; +bool g_EnhancedErrorChecking = false; const struct _cl_icd_dispatch* g_pNextDispatch = NULL; @@ -285,6 +284,8 @@ CL_API_ENTRY cl_int CL_API_CALL clInitLayer( _init_dispatch(); + getControl("CMDBUFEMU_EnhancedErrorChecking", g_EnhancedErrorChecking); + g_pNextDispatch = target_dispatch; *layer_dispatch_ret = &dispatch; From 916be459eb7fe651df9cea1d5248d059a9979f04 Mon Sep 17 00:00:00 2001 From: Ben Ashbaugh Date: Wed, 20 Nov 2024 21:36:20 -0800 Subject: [PATCH 4/4] add checks for a few more command buffer errors --- layers/10_cmdbufemu/emulate.cpp | 35 +++++++++++++++++++++++++++++++-- 1 file changed, 33 insertions(+), 2 deletions(-) diff --git a/layers/10_cmdbufemu/emulate.cpp b/layers/10_cmdbufemu/emulate.cpp index 64d672c9..6a1f2abd 100644 --- a/layers/10_cmdbufemu/emulate.cpp +++ b/layers/10_cmdbufemu/emulate.cpp @@ -1462,10 +1462,41 @@ typedef struct _cl_command_buffer_khr { return CL_INVALID_VALUE; } + if( ( event_wait_list == NULL && num_events_in_wait_list > 0 ) || + ( event_wait_list != NULL && num_events_in_wait_list == 0 ) ) + { + return CL_INVALID_EVENT_WAIT_LIST; + } + + cl_context cmdbuf_context = nullptr; + g_pNextDispatch->clGetCommandQueueInfo( + getQueue(), + CL_QUEUE_CONTEXT, + sizeof(cmdbuf_context), + &cmdbuf_context, + nullptr); + + for( cl_uint q = 0; q < num_queues && queues; q++ ) + { + if( queues[q] == NULL ) + { + return CL_INVALID_COMMAND_QUEUE; + } + + cl_context queue_context = nullptr; + g_pNextDispatch->clGetCommandQueueInfo( + queues[q], + CL_QUEUE_CONTEXT, + sizeof(queue_context), + &queue_context, + nullptr); + if( queue_context != cmdbuf_context ) + { + return CL_INVALID_CONTEXT; + } + } // CL_INCOMPATIBLE_COMMAND_QUEUE_KHR if any element of queues is not compatible with the command-queue set on command_buffer creation at the same list index. - // CL_INVALID_CONTEXT if any element of queues does not have the same context as the command-queue set on command_buffer creation at the same list indes. - // CL_INVALID_CONTEXT if the context associated with the command buffer and events in event_wait_list are not the same. return CL_SUCCESS; }