diff --git a/sycl/test-e2e/ProfilingTag/released_tag_event.cpp b/sycl/test-e2e/ProfilingTag/released_tag_event.cpp new file mode 100644 index 0000000000000..d826131441976 --- /dev/null +++ b/sycl/test-e2e/ProfilingTag/released_tag_event.cpp @@ -0,0 +1,56 @@ +// REQUIRES: aspect-ext_oneapi_queue_profiling_tag +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests that a profiling tag event still reports its timestamp when the +// previous tag's event was destroyed before this tag was submitted. Releasing +// the previous event lets a backend recycle its internal event object for the +// new tag while the released tag's timestamp write is still in flight; that +// write must not disturb the state of the new tag. + +#include "common.hpp" + +#include +#include + +constexpr size_t NumIters = 200; +constexpr size_t NumTagsPerIter = 10; + +int main() { + try { + sycl::queue Queue{sycl::property::queue::in_order()}; + + int Failures = 0; + std::vector KeptEvents; + KeptEvents.reserve(NumTagsPerIter); + + for (size_t I = 0; I < NumIters; ++I) { + KeptEvents.clear(); + for (size_t J = 0; J < NumTagsPerIter; ++J) { + // Submit a tag and drop its event straight away, without waiting for + // it. + auto DroppedEvent = std::make_unique( + sycl::ext::oneapi::experimental::submit_profiling_tag(Queue)); + DroppedEvent.reset(); + + KeptEvents.push_back( + sycl::ext::oneapi::experimental::submit_profiling_tag(Queue)); + } + + Queue.wait_and_throw(); + + for (sycl::event &E : KeptEvents) { + uint64_t End = + E.get_profiling_info(); + CHECK(Failures, End != 0) + } + } + + return Failures; + } catch (const sycl::exception &E) { + // A tag whose state was disturbed reports its timestamp as unavailable + // rather than as zero, so the failure arrives as an exception. + std::cout << "Unexpected exception: " << E.what() << std::endl; + return 1; + } +} diff --git a/unified-runtime/source/adapters/level_zero/v2/event.hpp b/unified-runtime/source/adapters/level_zero/v2/event.hpp index 208371062fda7..07337c57c817a 100644 --- a/unified-runtime/source/adapters/level_zero/v2/event.hpp +++ b/unified-runtime/source/adapters/level_zero/v2/event.hpp @@ -44,8 +44,15 @@ struct event_profiling_data_t { private: ze_event_handle_t hZeEvent; - uint64_t recordEventEndTimestamp = 0; - uint64_t adjustedEventEndTimestamp = 0; + // TODO: drop both alignas once GSD-13449 is fixed in the driver. + // On some platforms the device's 8-byte write of the recorded timestamp also + // rewrites the rest of the aligned 64-byte block around it, with stale + // contents. Host writes to members sharing that block are lost - including + // timestampRecorded, which then makes a recorded event report + // UR_RESULT_ERROR_PROFILING_INFO_NOT_AVAILABLE. Keep the timestamp in a block + // of its own; both alignas are needed for that. + alignas(64) uint64_t recordEventEndTimestamp = 0; + alignas(64) uint64_t adjustedEventEndTimestamp = 0; // Timer resolution in nanoseconds (converted from cycles/sec) double zeTimerResolution = 0;