[SYCL][XPTI] Add accessor events

sycl/include/CL/sycl/buffer.hpp

@@ -331,39 +331,46 @@ class buffer {
   template <access::mode Mode, access::target Target = access::target::device>
   accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
            ext::oneapi::accessor_property_list<>>
-  get_access(handler &CommandGroupHandler) {
+  get_access(
+      handler &CommandGroupHandler,
+      const detail::code_location CodeLoc = detail::code_location::current()) {
     return accessor<T, dimensions, Mode, Target, access::placeholder::false_t,
-                    ext::oneapi::accessor_property_list<>>(*this,
-                                                           CommandGroupHandler);
+                    ext::oneapi::accessor_property_list<>>(
+        *this, CommandGroupHandler, {}, CodeLoc);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
            access::placeholder::false_t, ext::oneapi::accessor_property_list<>>
-  get_access() {
+  get_access(
+      const detail::code_location CodeLoc = detail::code_location::current()) {
     return accessor<T, dimensions, mode, access::target::host_buffer,
                     access::placeholder::false_t,
-                    ext::oneapi::accessor_property_list<>>(*this);
+                    ext::oneapi::accessor_property_list<>>(*this, {}, CodeLoc);
   }
 
   template <access::mode mode, access::target target = access::target::device>
   accessor<T, dimensions, mode, target, access::placeholder::false_t,
            ext::oneapi::accessor_property_list<>>
-  get_access(handler &commandGroupHandler, range<dimensions> accessRange,
-             id<dimensions> accessOffset = {}) {
+  get_access(
+      handler &commandGroupHandler, range<dimensions> accessRange,
+      id<dimensions> accessOffset = {},
+      const detail::code_location CodeLoc = detail::code_location::current()) {
     return accessor<T, dimensions, mode, target, access::placeholder::false_t,
                     ext::oneapi::accessor_property_list<>>(
-        *this, commandGroupHandler, accessRange, accessOffset);
+        *this, commandGroupHandler, accessRange, accessOffset, {}, CodeLoc);
   }
 
   template <access::mode mode>
   accessor<T, dimensions, mode, access::target::host_buffer,
            access::placeholder::false_t, ext::oneapi::accessor_property_list<>>
-  get_access(range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
+  get_access(
+      range<dimensions> accessRange, id<dimensions> accessOffset = {},
+      const detail::code_location CodeLoc = detail::code_location::current()) {
     return accessor<T, dimensions, mode, access::target::host_buffer,
                     access::placeholder::false_t,
-                    ext::oneapi::accessor_property_list<>>(*this, accessRange,
-                                                           accessOffset);
+                    ext::oneapi::accessor_property_list<>>(
+        *this, accessRange, accessOffset, {}, CodeLoc);
   }
 
 #if __cplusplus > 201402L

sycl/source/detail/accessor_impl.cpp

@@ -10,6 +10,7 @@
 #include <CL/sycl/detail/buffer_impl.hpp>
 #include <detail/event_impl.hpp>
 #include <detail/scheduler/scheduler.hpp>
+#include <detail/xpti_registry.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -37,6 +38,14 @@ void addHostAccessorAndWait(Requirement *Req) {
       detail::Scheduler::getInstance().addHostAccessor(Req);
   Event->wait(Event);
 }
+
+void constructorNotification(void *BufferObj, void *AccessorObj,
+                             cl::sycl::access::target Target,
+                             cl::sycl::access::mode Mode,
+                             const detail::code_location &CodeLoc) {
+  XPTIRegistry::bufferAccessorNotification(
+      BufferObj, AccessorObj, (uint32_t)Target, (uint32_t)Mode, CodeLoc);
 }
-}
-}
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/buffer_impl.cpp

@@ -38,7 +38,6 @@ void buffer_impl::constructorNotification(const detail::code_location &CodeLoc,
 void buffer_impl::destructorNotification(void *UserObj) {
   XPTIRegistry::bufferDestructorNotification(UserObj);
 }
-
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/xpti_registry.cpp

@@ -16,44 +16,50 @@
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
-void XPTIRegistry::bufferConstructorNotification(
-    void *UserObj, const detail::code_location &CodeLoc) {
-  (void)CodeLoc;
-#ifdef XPTI_ENABLE_INSTRUMENTATION
-  GlobalHandler::instance().getXPTIRegistry().initializeFrameworkOnce();
-  if (!xptiTraceEnabled())
-    return;
-
-  uint64_t IId = xptiGetUniqueId();
+xpti::trace_event_data_t *
+XPTIRegistry::createTraceEvent(void *Obj, const char *ObjName, uint64_t &IId,
+                               const detail::code_location &CodeLoc,
+                               uint16_t TraceEventType) {
   std::string Name;
   if (CodeLoc.fileName()) {
     Name = std::string(CodeLoc.fileName()) + ":" +
            std::to_string(CodeLoc.lineNumber()) + ":" +
            std::to_string(CodeLoc.columnNumber());
   } else {
-    // We try to create a unique string for the buffer constructor call by
-    // combining it with the the created object address
     xpti::utils::StringHelper NG;
-    Name = NG.nameWithAddress<void *>("buffer", UserObj);
+    Name = NG.nameWithAddress<void *>(ObjName, Obj);
   }
-  xpti::offload_buffer_data_t BufConstr{(uintptr_t)UserObj};
-
   xpti::payload_t Payload(
       Name.c_str(), (CodeLoc.fileName() ? CodeLoc.fileName() : ""),
-      CodeLoc.lineNumber(), CodeLoc.columnNumber(), (void *)UserObj);
+      CodeLoc.lineNumber(), CodeLoc.columnNumber(), (void *)Obj);
 
-  // Constructor calls could be at different user-code locations; We create a
-  // new event based on the code location info and if this has been seen
-  // before, a previously created event will be returned.
-  xpti::trace_event_data_t *TraceEvent =
-      xptiMakeEvent(Name.c_str(), &Payload, xpti::trace_offload_buffer_event,
-                    xpti_at::active, &IId);
+  // Calls could be at different user-code locations; We create a new event
+  // based on the code location info and if this has been seen before, a
+  // previously created event will be returned.
+  return xptiMakeEvent(Name.c_str(), &Payload, TraceEventType, xpti_at::active,
+                       &IId);
+}
+void XPTIRegistry::bufferConstructorNotification(
+    void *UserObj, const detail::code_location &CodeLoc) {
+  (void)UserObj;
+  (void)CodeLoc;
+#ifdef XPTI_ENABLE_INSTRUMENTATION
+  GlobalHandler::instance().getXPTIRegistry().initializeFrameworkOnce();
+  if (!xptiTraceEnabled())
+    return;
+
+  uint64_t IId;
+  xpti::offload_buffer_data_t BufConstr{(uintptr_t)UserObj};
+
+  xpti::trace_event_data_t *TraceEvent = createTraceEvent(
+      UserObj, "buffer", IId, CodeLoc, xpti::trace_offload_buffer_event);
   xptiNotifySubscribers(GBufferStreamID, xpti::trace_offload_alloc_construct,
                         nullptr, TraceEvent, IId, &BufConstr);
 #endif
 }
 
 void XPTIRegistry::bufferAssociateNotification(void *UserObj, void *MemObj) {
+  (void)UserObj;
   (void)MemObj;
 #ifdef XPTI_ENABLE_INSTRUMENTATION
   if (!xptiTraceEnabled())
@@ -69,6 +75,7 @@ void XPTIRegistry::bufferAssociateNotification(void *UserObj, void *MemObj) {
 }
 
 void XPTIRegistry::bufferReleaseNotification(void *UserObj, void *MemObj) {
+  (void)UserObj;
   (void)MemObj;
 #ifdef XPTI_ENABLE_INSTRUMENTATION
   if (!xptiTraceEnabled())
@@ -84,6 +91,7 @@ void XPTIRegistry::bufferReleaseNotification(void *UserObj, void *MemObj) {
 }
 
 void XPTIRegistry::bufferDestructorNotification(void *UserObj) {
+  (void)UserObj;
 #ifdef XPTI_ENABLE_INSTRUMENTATION
   if (!xptiTraceEnabled())
     return;
@@ -95,6 +103,29 @@ void XPTIRegistry::bufferDestructorNotification(void *UserObj) {
 #endif
 }
 
+void XPTIRegistry::bufferAccessorNotification(
+    void *UserObj, void *AccessorObj, uint32_t Target, uint32_t Mode,
+    const detail::code_location &CodeLoc) {
+  (void)UserObj;
+  (void)AccessorObj;
+  (void)CodeLoc;
+  (void)Target;
+  (void)Mode;
+#ifdef XPTI_ENABLE_INSTRUMENTATION
+  if (!xptiTraceEnabled())
+    return;
+
+  uint64_t IId;
+  xpti::offload_accessor_data_t AccessorConstr{
+      (uintptr_t)UserObj, (uintptr_t)AccessorObj, Target, Mode};
+
+  xpti::trace_event_data_t *TraceEvent = createTraceEvent(
+      UserObj, "accessor", IId, CodeLoc, xpti::trace_offload_accessor_event);
+  xptiNotifySubscribers(GBufferStreamID, xpti::trace_offload_alloc_accessor,
+                        nullptr, TraceEvent, IId, &AccessorConstr);
+#endif
+}
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/xpti_registry.hpp

@@ -95,6 +95,13 @@ class XPTIRegistry {
   static void bufferAssociateNotification(void *UserObj, void *MemObj);
   static void bufferReleaseNotification(void *UserObj, void *MemObj);
   static void bufferDestructorNotification(void *UserObj);
+  static void bufferAccessorNotification(void *UserObj, void *AccessorObj,
+                                         uint32_t Target, uint32_t Mode,
+                                         const detail::code_location &CodeLoc);
+  static xpti::trace_event_data_t *
+  createTraceEvent(void *Obj, const char *ObjName, uint64_t &IId,
+                   const detail::code_location &CodeLoc,
+                   uint16_t TraceEventType);
 
 private:
   std::unordered_set<std::string> MActiveStreams;

sycl/test/abi/sycl_symbols_linux.dump

@@ -3882,6 +3882,7 @@ _ZN2cl4sycl6detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6devic
 _ZN2cl4sycl6detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKS5_INS0_9kernel_idESaISB_EENS0_12bundle_stateE
 _ZN2cl4sycl6detail22has_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE
 _ZN2cl4sycl6detail22has_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKS5_INS0_9kernel_idESaISB_EENS0_12bundle_stateE
+_ZN2cl4sycl6detail23constructorNotificationEPvS2_NS0_6access6targetENS3_4modeERKNS1_13code_locationE
 _ZN2cl4sycl6detail23getESIMDDeviceInterfaceEv
 _ZN2cl4sycl6detail24find_device_intersectionERKSt6vectorINS0_13kernel_bundleILNS0_12bundle_stateE1EEESaIS5_EE
 _ZN2cl4sycl6detail27getPixelCoordLinearFiltModeENS0_3vecIfLi4EEENS0_15addressing_modeENS0_5rangeILi3EEERS3_

sycl/test/abi/sycl_symbols_windows.dump

@@ -49,6 +49,7 @@
 ??$get_info@$0BABB@@device@sycl@cl@@QEBA_KXZ
 ??$get_info@$0BABBA@@device@sycl@cl@@QEBA_NXZ
 ??$get_info@$0BABBB@@device@sycl@cl@@QEBA?AV?$vector@W4memory_order@sycl@cl@@V?$allocator@W4memory_order@sycl@cl@@@std@@@std@@XZ
+??$get_info@$0BABBC@@device@sycl@cl@@QEBAIXZ
 ??$get_info@$0BABBC@@kernel@sycl@cl@@QEBAIAEBVdevice@12@@Z
 ??$get_info@$0BABC@@device@sycl@cl@@QEBA_KXZ
 ??$get_info@$0BABD@@device@sycl@cl@@QEBA_KXZ
@@ -148,7 +149,6 @@
 ??$get_info@$0EBJD@@device@sycl@cl@@QEBA_NXZ
 ??$get_info@$0EBJE@@device@sycl@cl@@QEBA_NXZ
 ??$get_info@$0ECBD@@device@sycl@cl@@QEBA_NXZ
-??$get_info@$0BABBC@@device@sycl@cl@@QEBAIXZ
 ??$get_info@$0JAA@@platform@sycl@cl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ
 ??$get_info@$0JAB@@platform@sycl@cl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ
 ??$get_info@$0JAC@@platform@sycl@cl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ
@@ -1457,6 +1457,7 @@
 ?compile_with_kernel_name@program@sycl@cl@@AEAAXV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@0_J@Z
 ?compile_with_source@program@sycl@cl@@QEAAXV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@0@Z
 ?constructorNotification@buffer_impl@detail@sycl@cl@@QEAAXAEBUcode_location@234@PEAX@Z
+?constructorNotification@detail@sycl@cl@@YAXPEAX0W4target@access@23@W4mode@523@AEBUcode_location@123@@Z
 ?contains_specialization_constants@kernel_bundle_plain@detail@sycl@cl@@QEBA_NXZ
 ?contextSetExtendedDeleter@pi@detail@sycl@cl@@YAXAEBVcontext@34@P6AXPEAX@Z1@Z
 ?convertChannelOrder@detail@sycl@cl@@YA?AW4_pi_image_channel_order@@W4image_channel_order@23@@Z

xpti/include/xpti/xpti_data_types.h

@@ -387,6 +387,8 @@ enum class trace_point_type_t : uint16_t {
   offload_alloc_destruct = XPTI_TRACE_POINT_BEGIN(22),
   /// Used to notify about releasing internal handle for offload buffer
   offload_alloc_release = XPTI_TRACE_POINT_BEGIN(23),
+  /// Used to notify about creation accessor for ofload buffer
+  offload_alloc_accessor = XPTI_TRACE_POINT_BEGIN(24),
   /// Indicates that the trace point is user defined and only the tool defined
   /// for a stream will be able to handle it
   user_defined = 1 << 7
@@ -451,6 +453,8 @@ enum class trace_event_type_t : uint16_t {
   offload_write = XPTI_EVENT(8),
   /// Indicates that the current event is an offload buffer related
   offload_buffer = XPTI_EVENT(9),
+  /// Indicates that the current event is an offload accessor related
+  offload_accessor = XPTI_EVENT(10),
   /// User defined event for extensibility and will have to be registered by
   /// the tool/runtime
   user_defined = 1 << 7
@@ -517,6 +521,19 @@ struct offload_buffer_data_t {
   /// A pointer to user level memory offload object.
   uintptr_t user_object_handle = 0;
 };
+
+/// Describes offload accessor
+struct offload_accessor_data_t {
+  /// A pointer to user level buffer offload object.
+  uintptr_t buffer_handle = 0;
+  /// A pointer to user level accessor offload object.
+  uintptr_t accessor_handle = 0;
+  /// Access target
+  uint32_t target = 0;
+  /// Access mode
+  uint32_t mode = 0;
+};
+
 /// Describes association between user level and platform specific
 /// offload buffer object
 struct offload_buffer_association_data_t {
@@ -632,13 +649,17 @@ constexpr uint16_t trace_offload_alloc_destruct =
     static_cast<uint16_t>(xpti::trace_point_type_t::offload_alloc_destruct);
 constexpr uint16_t trace_offload_alloc_release =
     static_cast<uint16_t>(xpti::trace_point_type_t::offload_alloc_release);
+constexpr uint16_t trace_offload_alloc_accessor =
+    static_cast<uint16_t>(xpti::trace_point_type_t::offload_alloc_accessor);
 
 constexpr uint16_t trace_graph_event =
     static_cast<uint16_t>(xpti::trace_event_type_t::graph);
 constexpr uint16_t trace_algorithm_event =
     static_cast<uint16_t>(xpti::trace_event_type_t::algorithm);
 constexpr uint16_t trace_offload_buffer_event =
     static_cast<uint16_t>(xpti::trace_event_type_t::offload_buffer);
+constexpr uint16_t trace_offload_accessor_event =
+    static_cast<uint16_t>(xpti::trace_event_type_t::offload_accessor);
 } // namespace xpti
 
 using xpti_tp = xpti::trace_point_type_t;

xptifw/src/xpti_trace_framework.cpp

@@ -440,6 +440,7 @@ class Tracepoints {
       // Add source file information ot string table
       source_id =
           MStringTableRef.add(Payload->source_file, &Payload->source_file);
+      line_no = Payload->line_no;
     }
     if ((Payload->flags &
          static_cast<uint64_t>(payload_flag_t::StackTraceAvailable))) {
@@ -448,15 +449,15 @@ class Tracepoints {
           MStringTableRef.add(Payload->stack_trace, &Payload->stack_trace);
     }
     // Pack the 1st 64-bit value with string ID from source file name and line
-    // number; pack the 2nd 54-bit value with stack backtrace string ID and the
+    // number; pack the 2nd 64-bit value with stack backtrace string ID and the
     // kernel name string ID
     Payload->uid.p1 = XPTI_PACK32_RET64(source_id, line_no);
     Payload->uid.p2 = XPTI_PACK32_RET64(stack_id, name_id);
     // The code pointer for the kernel is already in 64-bit format
     if ((Payload->flags &
          static_cast<uint64_t>(payload_flag_t::CodePointerAvailable)))
       Payload->uid.p3 = (uint64_t)Payload->code_ptr_va;
-    // Generate the had from the information available and this will be our
+    // Generate the hash from the information available and this will be our
     // unique ID for the trace point.
     HashValue = Payload->uid.hash();
     Payload->flags |= static_cast<uint64_t>(payload_flag_t::HashAvailable);
@@ -941,7 +942,7 @@ class Framework {
       // have 'nullptr' for both the Parent and Object only if UserData is
       // provided and the trace_point_type is function_begin/function_end.
       // This allows us to trace function calls without too much effort.
-      std::array<trace_point_type_t, 12> AllowedTypes = {
+      std::array<trace_point_type_t, 13> AllowedTypes = {
           trace_point_type_t::function_begin,
           trace_point_type_t::function_end,
           trace_point_type_t::function_with_args_begin,
@@ -953,7 +954,8 @@ class Framework {
           trace_point_type_t::offload_alloc_construct,
           trace_point_type_t::offload_alloc_associate,
           trace_point_type_t::offload_alloc_release,
-          trace_point_type_t::offload_alloc_destruct};
+          trace_point_type_t::offload_alloc_destruct,
+          trace_point_type_t::offload_alloc_accessor};
       const auto Predicate = [TraceType](trace_point_type_t RHS) {
         return TraceType == static_cast<uint16_t>(RHS);
       };

xptifw/unit_test/xpti_api_tests.cpp

@@ -275,6 +275,11 @@ TEST_F(xptiApiTest, xptiRegisterCallbackGoodInput) {
       fn_callback);
   EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_SUCCESS);
 
+  Result = xptiRegisterCallback(
+      StreamID, (uint16_t)xpti::trace_point_type_t::offload_alloc_accessor,
+      fn_callback);
+  EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_SUCCESS);
+
   Result = xptiRegisterCallback(
       StreamID, (uint16_t)xpti::trace_point_type_t::mem_alloc_begin,
       fn_callback);
@@ -355,6 +360,11 @@ TEST_F(xptiApiTest, xptiNotifySubscribersBadInput) {
       nullptr, nullptr, 0, nullptr);
   EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_INVALIDARG);
 
+  Result = xptiNotifySubscribers(
+      StreamID, (uint16_t)xpti::trace_point_type_t::offload_alloc_accessor,
+      nullptr, nullptr, 0, nullptr);
+  EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_INVALIDARG);
+
   Result = xptiNotifySubscribers(
       StreamID, (uint16_t)xpti::trace_point_type_t::mem_alloc_begin, nullptr,
       nullptr, 0, nullptr);
@@ -427,6 +437,7 @@ TEST_F(xptiApiTest, xptiNotifySubscribersGoodInput) {
   xpti::offload_buffer_data_t UserBufferData{0x01020304};
   xpti::offload_buffer_association_data_t AssociationData{0x01020304,
                                                           0x05060708};
+  xpti::offload_accessor_data_t UserAccessorData{0x01020304, 0x09000102, 1, 2};
 
   tmp = func_callback_update;
   Result = xptiNotifySubscribers(
@@ -445,6 +456,10 @@ TEST_F(xptiApiTest, xptiNotifySubscribersGoodInput) {
       StreamID, (uint16_t)xpti::trace_point_type_t::offload_alloc_destruct,
       nullptr, (xpti::trace_event_data_t *)1, 0, &UserBufferData);
   EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_SUCCESS);
+  Result = xptiNotifySubscribers(
+      StreamID, (uint16_t)xpti::trace_point_type_t::offload_alloc_accessor,
+      nullptr, (xpti::trace_event_data_t *)1, 0, &UserAccessorData);
+  EXPECT_EQ(Result, xpti::result_t::XPTI_RESULT_SUCCESS);
 
   Result = xptiRegisterCallback(
       StreamID, (uint16_t)xpti::trace_point_type_t::mem_alloc_begin,