diff --git a/include/getenv_util.hpp b/include/getenv_util.hpp new file mode 100644 index 0000000..9e24f92 --- /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 169b47e..418af38 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 532ecae..6a1f2ab 100644 --- a/layers/10_cmdbufemu/emulate.cpp +++ b/layers/10_cmdbufemu/emulate.cpp @@ -1221,9 +1221,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); } } @@ -1236,6 +1240,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 ) @@ -1281,7 +1298,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; } #if defined(cl_khr_command_buffer_mutable_dispatch) @@ -1432,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; } @@ -1476,6 +1537,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; } @@ -1576,9 +1654,84 @@ typedef struct _cl_command_buffer_khr #endif // defined(cl_khr_command_buffer_mutable_dispatch) std::atomic RefCount; + std::vector TestQueues; + std::vector BlockingEvents; + std::vector> Commands; std::atomic NextSyncPoint; + void setupTestQueue(cl_command_queue src) + { + if( g_EnhancedErrorChecking ) + { + 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 #if defined(cl_khr_command_buffer_mutable_dispatch) @@ -1947,6 +2100,23 @@ cl_int CL_API_CALL clCommandCopyBufferKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2004,6 +2174,26 @@ cl_int CL_API_CALL clCommandCopyBufferRectKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2061,6 +2251,22 @@ cl_int CL_API_CALL clCommandCopyBufferToImageKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2114,6 +2320,22 @@ cl_int CL_API_CALL clCommandCopyImageKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2167,6 +2389,22 @@ cl_int CL_API_CALL clCommandCopyImageToBufferKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2220,6 +2458,22 @@ cl_int CL_API_CALL clCommandFillBufferKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2272,6 +2526,21 @@ cl_int CL_API_CALL clCommandFillImageKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2322,6 +2591,21 @@ cl_int CL_API_CALL clCommandSVMMemcpyKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2372,6 +2656,21 @@ cl_int CL_API_CALL clCommandSVMMemFillKHR_EMU( { return CL_INVALID_VALUE; } + 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( @@ -2422,6 +2721,22 @@ cl_int CL_API_CALL clCommandNDRangeKernelKHR_EMU( return CL_INVALID_VALUE; } #endif // !defined(cl_khr_command_buffer_mutable_dispatch) + 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 70e653d..d0fa6a3 100644 --- a/layers/10_cmdbufemu/emulate.h +++ b/layers/10_cmdbufemu/emulate.h @@ -9,6 +9,10 @@ #include +extern bool g_EnhancedErrorChecking; + +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..1194b96 100644 --- a/layers/10_cmdbufemu/main.cpp +++ b/layers/10_cmdbufemu/main.cpp @@ -23,10 +23,17 @@ #include #include +#include "getenv_util.hpp" #include "layer_util.hpp" #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. + +bool g_EnhancedErrorChecking = false; + const struct _cl_icd_dispatch* g_pNextDispatch = NULL; static cl_int CL_API_CALL @@ -277,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;