[SYCL][XPTI] Enable PI calls notifications with arguments

sycl/include/CL/sycl/detail/common.hpp

@@ -31,6 +31,9 @@ namespace detail {
 constexpr const char *SYCL_STREAM_NAME = "sycl";
 // Stream name being used for traces generated from the SYCL plugin layer
 constexpr const char *SYCL_PICALL_STREAM_NAME = "sycl.pi";
+// Stream name being used for traces generated from PI calls. This stream
+// contains information about function arguments.
+constexpr const char *SYCL_PIARGCALL_STREAM_NAME = "sycl.pi.arg";
 // Data structure that captures the user code location information using the
 // builtin capabilities of the compiler
 struct code_location {

sycl/include/CL/sycl/detail/pi.hpp

@@ -182,6 +182,23 @@ uint64_t emitFunctionBeginTrace(const char *FName);
 /// \param FName The name of the PI API call
 void emitFunctionEndTrace(uint64_t CorrelationID, const char *FName);
 
+/// Notifies XPTI subscribers about PI function calls and packs call arguments.
+///
+/// \param FuncID is the API hash ID from PiApiID type trait.
+/// \param ArgsData is a pointer to packed function call arguments.
+uint64_t emitFunctionWithArgsBeginTrace(uint32_t FuncID,
+                                        unsigned char *ArgsData);
+
+/// Notifies XPTI subscribers about PI function call result.
+///
+/// \param CorrelationID The correlation ID for the API call generated by the
+/// emitFunctionWithArgsBeginTrace() call.
+/// \param FuncID is the API hash ID from PiApiID type trait.
+/// \param ArgsData is a pointer to packed function call arguments.
+/// \param Result is function call result value.
+void emitFunctionWithArgsEndTrace(uint64_t CorrelationID, uint32_t FuncID,
+                                  unsigned char *ArgsData, pi_result Result);
+
 // A wrapper for passing around byte array properties
 class ByteArray {
 public:
@@ -391,3 +408,5 @@ namespace RT = cl::sycl::detail::pi;
 
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
+
+#undef _PI_API

sycl/include/CL/sycl/detail/pi_api_id.hpp

@@ -0,0 +1,60 @@
+//==---------- pi_api_id.hpp - PI API function IDs -------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+/// \file pi_api_id.hpp
+/// This file contains mapping between PI API functions and their IDs. A hash
+/// function is used to generate IDs. The reason for usage of a hash function
+/// instead of raw values of PiApiKind is ABI stability. New functions can be
+/// added to PiApiKind, and there's no reliable way for external XPTI users to
+/// know when to update the values. Hashes are calculated from API function
+/// name, so it will remain the same after update to pi.def.
+///
+/// \ingroup sycl_pi
+
+#pragma once
+
+#include <CL/sycl/detail/pi.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+constexpr uint32_t cxpow(uint32_t Base, uint32_t Pow) {
+  uint32_t Res = Base;
+  for (uint32_t I = 1; I < Pow; I++)
+    Res *= Base;
+  return Res;
+}
+
+/// This is a simple implementation of polynomial rolling hash function.
+///
+/// The general formula for the hash is Sum(s[i] * p^i) mod m.
+/// Since only English characters are used for PI function names, p = 53 is
+/// chosen. m = 1051 is a fairly big prime number for the task.
+constexpr uint32_t cxhash(const char *Str) {
+  constexpr uint32_t p = 53;
+  constexpr uint32_t m = 1051;
+  uint32_t Hash = 0;
+  uint32_t Len = 0;
+  while (Str[Len++] != '\0')
+    Hash += Str[Len - 1] * cxpow(p, Len - 1);
+  return Hash % m;
+}
+
+template <PiApiKind Api> struct PiApiID {};
+
+#define _PI_API(api)                                                           \
+  template <> struct PiApiID<PiApiKind::api> {                                 \
+    static constexpr uint32_t id = cxhash(#api);                               \
+  };
+
+#include <CL/sycl/detail/pi.def>
+
+#undef _PI_API
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)