diff --git a/sycl/include/sycl/ext/oneapi/experimental/profiling_tag.hpp b/sycl/include/sycl/ext/oneapi/experimental/profiling_tag.hpp index dec91f77a7c7c..1be7585b7f272 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/profiling_tag.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/profiling_tag.hpp @@ -21,22 +21,24 @@ namespace ext::oneapi::experimental { inline event submit_profiling_tag(queue &Queue, const sycl::detail::code_location &CodeLoc = sycl::detail::code_location::current()) { - if (Queue.get_device().has(aspect::ext_oneapi_queue_profiling_tag)) { - return Queue.submit( - [=](handler &CGH) { - sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH); - }, - CodeLoc); - } - - // If it is not supported natively on the device, we use another path if - // profiling is enabled. - if (!Queue.has_property()) + // The profiling tag can be serviced natively when the device advertises the + // ext_oneapi_queue_profiling_tag aspect. Otherwise, we can still service it + // as long as the queue has profiling enabled. In both cases we submit an + // internal profiling-tag command group: the runtime records the timestamp + // using a native device command where possible and otherwise falls back to a + // barrier (see CGType::ProfilingTag handling in the scheduler). + if (!Queue.get_device().has(aspect::ext_oneapi_queue_profiling_tag) && + !Queue.has_property()) throw sycl::exception( make_error_code(errc::invalid), "Device must either have aspect::ext_oneapi_queue_profiling_tag or the " "queue must have profiling enabled."); - return Queue.ext_oneapi_submit_barrier(); + + return Queue.submit( + [=](handler &CGH) { + sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH); + }, + CodeLoc); } } // namespace ext::oneapi::experimental diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 1871151a79e5a..9f4941b36231e 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3651,17 +3651,13 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { assert(MQueue && "Profiling tag requires a valid queue"); adapter_impl &Adapter = MQueue->getAdapter(); - bool IsInOrderQueue = MQueue->isInOrder(); - ur_event_handle_t *TimestampDeps = nullptr; - size_t NumTimestampDeps = 0; - - // TO DO - once the following WA removed: to change call to call_nocheck and - // return operation result to Command::enqueue (see other CG types). Set - // UREvent to EventImpl only for successful case. + const bool IsInOrderQueue = MQueue->isInOrder(); // If the queue is not in-order, the implementation will need to first // insert a marker event that the timestamp waits for. ur_event_handle_t PreTimestampMarkerEvent{}; + ur_event_handle_t *TimestampDeps = nullptr; + size_t NumTimestampDeps = 0; if (!IsInOrderQueue) { // FIXME: urEnqueueEventsWait on the L0 adapter requires a double-release. // Use that instead once it has been fixed. @@ -3674,25 +3670,51 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { NumTimestampDeps = 1; } - Adapter.call( - MQueue->getHandleRef(), - /*blocking=*/false, NumTimestampDeps, TimestampDeps, Event); - - // If the queue is not in-order, we need to insert a barrier. This barrier - // does not need output events as it will implicitly enforce the following - // enqueue is blocked until it finishes. - if (!IsInOrderQueue) { - // We also need to release the timestamp event from the marker. - Adapter.call(PreTimestampMarkerEvent); - // FIXME: Due to a bug in the L0 UR adapter, we will leak events if we do - // not pass an output event to the UR call. Once that is fixed, - // this immediately-deleted event can be removed. - ur_event_handle_t PostTimestampBarrierEvent{}; - Adapter.call( - MQueue->getHandleRef(), - /*num_events_in_wait_list=*/0, - /*event_wait_list=*/nullptr, &PostTimestampBarrierEvent); - Adapter.call(PostTimestampBarrierEvent); + // Try to record a device timestamp natively. Not every backend can do so: + // the OpenCL backend can only record a reliable timestamp on a + // profiling-enabled queue (see intel/llvm#22229). When the recording is + // unsupported we fall back to a plain barrier, whose event still carries + // (best-effort) profiling information and provides the same ordering. + ur_result_t TimestampResult = + Adapter.call_nocheck( + MQueue->getHandleRef(), + /*blocking=*/false, NumTimestampDeps, TimestampDeps, Event); + + if (TimestampResult == UR_RESULT_ERROR_UNSUPPORTED_FEATURE) { + // Release the marker the (unused) timestamp recording would have waited + // for. + if (!IsInOrderQueue) + Adapter.call(PreTimestampMarkerEvent); + // A barrier with an empty wait list waits for all previously-submitted + // work and blocks subsequent work, providing the same ordering semantics + // as a native profiling tag. + if (auto Result = + Adapter.call_nocheck( + MQueue->getHandleRef(), + /*num_events_in_wait_list=*/0, + /*event_wait_list=*/nullptr, Event); + Result != UR_RESULT_SUCCESS) + return Result; + } else { + if (TimestampResult != UR_RESULT_SUCCESS) + return TimestampResult; + + // If the queue is not in-order, we need to insert a barrier. This barrier + // does not need output events as it will implicitly enforce the following + // enqueue is blocked until it finishes. + if (!IsInOrderQueue) { + // We also need to release the timestamp event from the marker. + Adapter.call(PreTimestampMarkerEvent); + // FIXME: Due to a bug in the L0 UR adapter, we will leak events if we + // do not pass an output event to the UR call. Once that is + // fixed, this immediately-deleted event can be removed. + ur_event_handle_t PostTimestampBarrierEvent{}; + Adapter.call( + MQueue->getHandleRef(), + /*num_events_in_wait_list=*/0, + /*event_wait_list=*/nullptr, &PostTimestampBarrierEvent); + Adapter.call(PostTimestampBarrierEvent); + } } SetEventHandleOrDiscard(); diff --git a/sycl/test-e2e/ProfilingTag/in_order_profiling_queue.cpp b/sycl/test-e2e/ProfilingTag/in_order_profiling_queue.cpp index fa08c8d41f96f..dcbcba83a9877 100644 --- a/sycl/test-e2e/ProfilingTag/in_order_profiling_queue.cpp +++ b/sycl/test-e2e/ProfilingTag/in_order_profiling_queue.cpp @@ -8,11 +8,6 @@ // Note: Extension should work even on devices that do not support the // ext_oneapi_queue_profiling_tag aspect. -// Bug in OpenCL GPU driver causes fallback solution to return end time later -// than the submission of the following work. -// UNSUPPORTED: opencl && gpu -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22229 - // HIP backend currently returns invalid values for submission time queries. // UNSUPPORTED: hip // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/12904 diff --git a/sycl/test-e2e/ProfilingTag/profiling_queue.cpp b/sycl/test-e2e/ProfilingTag/profiling_queue.cpp index f85739fedd92d..2431bb7c719b6 100644 --- a/sycl/test-e2e/ProfilingTag/profiling_queue.cpp +++ b/sycl/test-e2e/ProfilingTag/profiling_queue.cpp @@ -8,11 +8,6 @@ // Note: Extension should work even on devices that do not support the // ext_oneapi_queue_profiling_tag aspect. -// Bug in OpenCL GPU driver causes fallback solution to return end time later -// than the submission of the following work. -// UNSUPPORTED: opencl && gpu -// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/22229 - // HIP backend currently returns invalid values for submission time queries. // UNSUPPORTED: hip // UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/12904 diff --git a/sycl/unittests/Extensions/ProfilingTag.cpp b/sycl/unittests/Extensions/ProfilingTag.cpp index 924eaf6bd66ff..f6fc0a6838ea9 100644 --- a/sycl/unittests/Extensions/ProfilingTag.cpp +++ b/sycl/unittests/Extensions/ProfilingTag.cpp @@ -30,6 +30,14 @@ inline ur_result_t after_urEnqueueTimestampRecordingExp(void *) { return UR_RESULT_SUCCESS; } +// Simulates a backend that cannot record a device timestamp (e.g. the OpenCL +// backend on a queue without profiling enabled), so that the scheduler falls +// back to a barrier. +inline ur_result_t replace_urEnqueueTimestampRecordingExpUnsupported(void *) { + ++counter_urEnqueueTimestampRecordingExp; + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + thread_local std::optional LatestProfilingQuery; inline ur_result_t after_urEventGetProfilingInfo(void *pParams) { auto &Params = @@ -205,24 +213,57 @@ TEST_F(ProfilingTagTest, ProfilingTagFallbackDefaultQueue) { } } -TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueue) { +// Without the aspect but with profiling enabled, the tag is still serviced via +// the profiling-tag command group. If the backend can record a device +// timestamp (as the OpenCL backend now does on a profiling-enabled queue), the +// native recording path is used rather than a bare barrier. +TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueueTimestamp) { mock::getCallbacks().set_after_callback("urDeviceGetInfo", &after_urDeviceGetInfo); mock::getCallbacks().set_after_callback( "urEnqueueTimestampRecordingExp", &after_urEnqueueTimestampRecordingExp); mock::getCallbacks().set_after_callback( - "urEnqueueEventsWaitWithBarrierExt", + "urEnqueueEventsWaitWithBarrier", &after_urEnqueueEventsWaitWithBarrierExt); sycl::context Ctx{sycl::platform()}; sycl::queue Queue{Ctx, sycl::default_selector_v, - {sycl::property::queue::enable_profiling()}}; + {sycl::property::queue::enable_profiling(), + sycl::property::queue::in_order()}}; + sycl::device Dev = Queue.get_device(); + + ASSERT_FALSE(Dev.has(sycl::aspect::ext_oneapi_queue_profiling_tag)); + + sycl::event E = sycl::ext::oneapi::experimental::submit_profiling_tag(Queue); + ASSERT_EQ(size_t{1}, counter_urEnqueueTimestampRecordingExp); + ASSERT_EQ(size_t{0}, counter_urEnqueueEventsWaitWithBarrierExt); +} + +// If the backend reports that it cannot record a device timestamp, the +// scheduler must gracefully fall back to a barrier. +TEST_F(ProfilingTagTest, ProfilingTagFallbackProfilingQueueBarrier) { + mock::getCallbacks().set_after_callback("urDeviceGetInfo", + &after_urDeviceGetInfo); + mock::getCallbacks().set_replace_callback( + "urEnqueueTimestampRecordingExp", + &replace_urEnqueueTimestampRecordingExpUnsupported); + mock::getCallbacks().set_after_callback( + "urEnqueueEventsWaitWithBarrier", + &after_urEnqueueEventsWaitWithBarrierExt); + + sycl::context Ctx{sycl::platform()}; + sycl::queue Queue{Ctx, + sycl::default_selector_v, + {sycl::property::queue::enable_profiling(), + sycl::property::queue::in_order()}}; sycl::device Dev = Queue.get_device(); ASSERT_FALSE(Dev.has(sycl::aspect::ext_oneapi_queue_profiling_tag)); sycl::event E = sycl::ext::oneapi::experimental::submit_profiling_tag(Queue); - ASSERT_EQ(size_t{0}, counter_urEnqueueTimestampRecordingExp); + // The timestamp recording is attempted and reports unsupported, so a single + // barrier is used as the fallback on this in-order queue. + ASSERT_EQ(size_t{1}, counter_urEnqueueTimestampRecordingExp); ASSERT_EQ(size_t{1}, counter_urEnqueueEventsWaitWithBarrierExt); } diff --git a/unified-runtime/source/adapters/opencl/context.hpp b/unified-runtime/source/adapters/opencl/context.hpp index 0428bc944edb5..fb88d371496a7 100644 --- a/unified-runtime/source/adapters/opencl/context.hpp +++ b/unified-runtime/source/adapters/opencl/context.hpp @@ -13,6 +13,7 @@ #include "common/ur_ref_count.hpp" #include "device.hpp" +#include #include namespace ur::opencl { @@ -25,6 +26,15 @@ struct ur_context_handle_t_ : handle_base { bool IsNativeHandleOwned = true; ur::RefCount RefCount; + // Small device buffer used to emulate a timestamp recording via a real, + // profilable command (see urEnqueueTimestampRecordingExp). Some OpenCL + // drivers timestamp synchronization-only commands (barriers/markers) at the + // point they are inserted into the pipeline rather than when the preceding + // work actually completes, which breaks profiling-tag orderings + // (see https://github.com/intel/llvm/issues/22229). Created lazily. + std::mutex TimestampRecordingBufferMutex; + cl_mem TimestampRecordingBuffer = nullptr; + ur_context_handle_t_(const ur_context_handle_t_ &) = delete; ur_context_handle_t_ &operator=(const ur_context_handle_t_ &) = delete; @@ -40,6 +50,21 @@ struct ur_context_handle_t_ : handle_base { static ur_result_t makeWithNative(native_type Ctx, uint32_t DevCount, const ur_device_handle_t *phDevices, ur_context_handle_t &Context); + + // Returns the small internal buffer used by urEnqueueTimestampRecordingExp, + // creating it on first use. Thread-safe. + ur_result_t getTimestampRecordingBuffer(cl_mem *OutBuffer) { + std::lock_guard Lock(TimestampRecordingBufferMutex); + if (!TimestampRecordingBuffer) { + cl_int CLErr = CL_SUCCESS; + TimestampRecordingBuffer = clCreateBuffer( + CLContext, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &CLErr); + CL_RETURN_ON_FAILURE(CLErr); + } + *OutBuffer = TimestampRecordingBuffer; + return UR_RESULT_SUCCESS; + } + ~ur_context_handle_t_() noexcept { // If we're reasonably sure this context is about to be destroyed we should // clear the ext function pointer cache. This isn't foolproof sadly but it @@ -50,6 +75,9 @@ struct ur_context_handle_t_ : handle_base { for (uint32_t i = 0; i < DeviceCount; i++) { ur::opencl::urDeviceRelease(cast(Devices[i])); } + if (TimestampRecordingBuffer) { + clReleaseMemObject(TimestampRecordingBuffer); + } if (IsNativeHandleOwned) { clReleaseContext(CLContext); } diff --git a/unified-runtime/source/adapters/opencl/event.cpp b/unified-runtime/source/adapters/opencl/event.cpp index 762d54009e3d6..ad887d88c3f39 100644 --- a/unified-runtime/source/adapters/opencl/event.cpp +++ b/unified-runtime/source/adapters/opencl/event.cpp @@ -309,10 +309,52 @@ urEventSetCallback(ur_event_handle_t hEvent, ur_execution_info_t execStatus, return UR_RESULT_SUCCESS; } -UR_APIEXPORT ur_result_t UR_APICALL -urEnqueueTimestampRecordingExp(ur_queue_handle_t, bool, uint32_t, - const ur_event_handle_t *, ur_event_handle_t *) { - return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +UR_APIEXPORT ur_result_t UR_APICALL urEnqueueTimestampRecordingExp( + ur_queue_handle_t hQueue, bool Blocking, uint32_t numEventsInWaitList, + const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) { + // OpenCL has no native "record device timestamp" primitive. Some drivers + // (notably the Intel GPU driver) timestamp synchronization-only commands + // (barriers/markers) at the point they are inserted into the pipeline rather + // than when the preceding work actually completes, which breaks profiling-tag + // orderings (see https://github.com/intel/llvm/issues/22229). We therefore + // emulate a timestamp by enqueuing a lightweight *real* device command (a + // small buffer fill), whose profiling timestamps do reflect completion of the + // preceding work. + // + // This relies on OpenCL command profiling, which is only available when the + // queue was created with CL_QUEUE_PROFILING_ENABLE. If profiling is not + // enabled we cannot honor the request and report it as unsupported so the + // caller can fall back. + auto Queue = cast(hQueue); + + cl_command_queue_properties QueueProperties = 0; + CL_RETURN_ON_FAILURE(clGetCommandQueueInfo( + Queue->CLQueue, CL_QUEUE_PROPERTIES, sizeof(QueueProperties), + &QueueProperties, nullptr)); + if (!(QueueProperties & CL_QUEUE_PROFILING_ENABLE)) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + cl_mem Buffer = nullptr; + UR_RETURN_ON_FAILURE(Queue->Context->getTimestampRecordingBuffer(&Buffer)); + + std::vector CLWaitEvents(numEventsInWaitList); + for (uint32_t I = 0; I < numEventsInWaitList; I++) + CLWaitEvents[I] = cast(phEventWaitList[I])->CLEvent; + + const cl_uint Pattern = 0; + cl_event Event = nullptr; + CL_RETURN_ON_FAILURE(clEnqueueFillBuffer( + Queue->CLQueue, Buffer, &Pattern, sizeof(Pattern), /*offset=*/0, + /*size=*/sizeof(Pattern), numEventsInWaitList, CLWaitEvents.data(), + ifUrEvent(phEvent, Event))); + + UR_RETURN_ON_FAILURE( + createUREvent(Event, cast(Queue->Context), cast(Queue), phEvent)); + + if (Blocking) + CL_RETURN_ON_FAILURE(clFinish(Queue->CLQueue)); + + return UR_RESULT_SUCCESS; } UR_APIEXPORT ur_result_t UR_APICALL