[SYCL] Instrumentation for application profiling

buildbot/configure.py

@@ -10,10 +10,11 @@ def do_configure(args):
     llvm_dir = os.path.join(args.src_dir, "llvm")
     sycl_dir = os.path.join(args.src_dir, "sycl")
     spirv_dir = os.path.join(args.src_dir, "llvm-spirv")
+    xpti_dir = os.path.join(args.src_dir, "xpti")
     ocl_header_dir = os.path.join(args.obj_dir, "OpenCL-Headers")
     icd_loader_lib = os.path.join(args.obj_dir, "OpenCL-ICD-Loader", "build")
     llvm_targets_to_build = 'X86'
-    llvm_enable_projects = 'clang;llvm-spirv;sycl;opencl-aot'
+    llvm_enable_projects = 'clang;llvm-spirv;sycl;opencl-aot;xpti'
     libclc_targets_to_build = ''
     sycl_build_pi_cuda = 'OFF'
     llvm_enable_assertions = 'ON'
@@ -44,9 +45,10 @@ def do_configure(args):
         "-DCMAKE_BUILD_TYPE={}".format(args.build_type),
         "-DLLVM_ENABLE_ASSERTIONS={}".format(llvm_enable_assertions),
         "-DLLVM_TARGETS_TO_BUILD={}".format(llvm_targets_to_build),
-        "-DLLVM_EXTERNAL_PROJECTS=sycl;llvm-spirv;opencl-aot",
+        "-DLLVM_EXTERNAL_PROJECTS=sycl;llvm-spirv;opencl-aot;xpti",
         "-DLLVM_EXTERNAL_SYCL_SOURCE_DIR={}".format(sycl_dir),
         "-DLLVM_EXTERNAL_LLVM_SPIRV_SOURCE_DIR={}".format(spirv_dir),
+        "-DLLVM_EXTERNAL_XPTI_SOURCE_DIR={}".format(xpti_dir),
         "-DLLVM_ENABLE_PROJECTS={}".format(llvm_enable_projects),
         "-DLIBCLC_TARGETS_TO_BUILD={}".format(libclc_targets_to_build),
         "-DOpenCL_INCLUDE_DIR={}".format(ocl_header_dir),
@@ -57,6 +59,7 @@ def do_configure(args):
         "-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
         "-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.
         "-DLLVM_ENABLE_DOXYGEN={}".format(llvm_enable_doxygen),
+        "-DSYCL_ENABLE_XPTI_TRACING=ON", # Explicitly turn on XPTI tracing
         llvm_dir
     ]
 

sycl/CMakeLists.txt

@@ -26,6 +26,10 @@ if(SYCL_ENABLE_WERROR)
   endif()
 endif()
 
+# Create a soft option for enabling or disabling the instrumentation
+# of the SYCL runtime and expect enabling
+option(SYCL_ENABLE_XPTI_TRACING "Enable tracing of SYCL constructs" OFF)
+
 if(MSVC)
   set_property(GLOBAL PROPERTY USE_FOLDERS ON)
   # Skip asynchronous C++ exceptions catching and assume "extern C" functions
@@ -218,6 +222,13 @@ add_custom_target( sycl-toolchain
   COMMENT "Building SYCL compiler toolchain..."
 )
 
+if (SYCL_ENABLE_XPTI_TRACING)
+  add_dependencies( sycl-toolchain xpti)
+  if (MSVC)
+    add_dependencies( sycl-toolchain xptid)
+  endif()
+endif()
+
 if (NOT DEFINED LLVM_INCLUDE_TESTS)
   set(LLVM_INCLUDE_TESTS ON)
 endif()

sycl/include/CL/sycl/detail/cg.hpp

@@ -368,11 +368,21 @@ class CG {
      vector_class<detail::AccessorImplPtr> AccStorage,
      vector_class<shared_ptr_class<const void>> SharedPtrStorage,
      vector_class<Requirement *> Requirements,
-     vector_class<detail::EventImplPtr> Events)
+     vector_class<detail::EventImplPtr> Events, detail::code_location loc = {})
       : MType(Type), MArgsStorage(std::move(ArgsStorage)),
         MAccStorage(std::move(AccStorage)),
         MSharedPtrStorage(std::move(SharedPtrStorage)),
-        MRequirements(std::move(Requirements)), MEvents(std::move(Events)) {}
+        MRequirements(std::move(Requirements)), MEvents(std::move(Events)) {
+    // Capture the user code-location from Q.submit(), Q.parallel_for()
+    // etc for later use; if code location information is not available,
+    // the file name and function name members will be empty strings
+    if (loc.functionName())
+      MFunctionName = loc.functionName();
+    if (loc.fileName())
+      MFileName = loc.fileName();
+    MLine = loc.lineNumber();
+    MColumn = loc.columnNumber();
+  }
 
   CG(CG &&CommandGroup) = default;
 
@@ -397,6 +407,12 @@ class CG {
   vector_class<Requirement *> MRequirements;
   // List of events that order the execution of this CG
   vector_class<detail::EventImplPtr> MEvents;
+  // Member variables to capture the user code-location
+  // information from Q.submit(), Q.parallel_for() etc
+  // Storage for function name and source file name
+  string_class MFunctionName, MFileName;
+  // Storage for line and column of code location
+  int32_t MLine, MColumn;
 };
 
 // The class which represents "execute kernel" command group.
@@ -420,10 +436,10 @@ class CGExecKernel : public CG {
                vector_class<ArgDesc> Args, string_class KernelName,
                detail::OSModuleHandle OSModuleHandle,
                vector_class<shared_ptr_class<detail::stream_impl>> Streams,
-               CGTYPE Type)
+               CGTYPE Type, detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MNDRDesc(std::move(NDRDesc)), MHostKernel(std::move(HKernel)),
         MSyclKernel(std::move(SyclKernel)), MArgs(std::move(Args)),
         MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle),
@@ -450,10 +466,11 @@ class CGCopy : public CG {
          vector_class<detail::AccessorImplPtr> AccStorage,
          vector_class<shared_ptr_class<const void>> SharedPtrStorage,
          vector_class<Requirement *> Requirements,
-         vector_class<detail::EventImplPtr> Events)
+         vector_class<detail::EventImplPtr> Events,
+         detail::code_location loc = {})
       : CG(CopyType, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MSrc(Src), MDst(Dst) {}
   void *getSrc() { return MSrc; }
   void *getDst() { return MDst; }
@@ -470,10 +487,11 @@ class CGFill : public CG {
          vector_class<detail::AccessorImplPtr> AccStorage,
          vector_class<shared_ptr_class<const void>> SharedPtrStorage,
          vector_class<Requirement *> Requirements,
-         vector_class<detail::EventImplPtr> Events)
+         vector_class<detail::EventImplPtr> Events,
+         detail::code_location loc = {})
       : CG(FILL, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MPattern(std::move(Pattern)), MPtr((Requirement *)Ptr) {}
   Requirement *getReqToFill() { return MPtr; }
 };
@@ -487,10 +505,11 @@ class CGUpdateHost : public CG {
                vector_class<detail::AccessorImplPtr> AccStorage,
                vector_class<shared_ptr_class<const void>> SharedPtrStorage,
                vector_class<Requirement *> Requirements,
-               vector_class<detail::EventImplPtr> Events)
+               vector_class<detail::EventImplPtr> Events,
+               detail::code_location loc = {})
       : CG(UPDATE_HOST, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MPtr((Requirement *)Ptr) {}
 
   Requirement *getReqToUpdate() { return MPtr; }
@@ -508,10 +527,11 @@ class CGCopyUSM : public CG {
             vector_class<detail::AccessorImplPtr> AccStorage,
             vector_class<shared_ptr_class<const void>> SharedPtrStorage,
             vector_class<Requirement *> Requirements,
-            vector_class<detail::EventImplPtr> Events)
+            vector_class<detail::EventImplPtr> Events,
+            detail::code_location loc = {})
       : CG(COPY_USM, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MSrc(Src), MDst(Dst), MLength(Length) {}
 
   void *getSrc() { return MSrc; }
@@ -531,10 +551,11 @@ class CGFillUSM : public CG {
             vector_class<detail::AccessorImplPtr> AccStorage,
             vector_class<shared_ptr_class<const void>> SharedPtrStorage,
             vector_class<Requirement *> Requirements,
-            vector_class<detail::EventImplPtr> Events)
+            vector_class<detail::EventImplPtr> Events,
+            detail::code_location loc = {})
       : CG(FILL_USM, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MPattern(std::move(Pattern)), MDst(DstPtr), MLength(Length) {}
   void *getDst() { return MDst; }
   size_t getLength() { return MLength; }
@@ -552,10 +573,11 @@ class CGPrefetchUSM : public CG {
                 vector_class<detail::AccessorImplPtr> AccStorage,
                 vector_class<shared_ptr_class<const void>> SharedPtrStorage,
                 vector_class<Requirement *> Requirements,
-                vector_class<detail::EventImplPtr> Events)
+                vector_class<detail::EventImplPtr> Events,
+                detail::code_location loc = {})
       : CG(PREFETCH_USM, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MDst(DstPtr), MLength(Length) {}
   void *getDst() { return MDst; }
   size_t getLength() { return MLength; }
@@ -570,10 +592,11 @@ class CGInteropTask : public CG {
                 std::vector<detail::AccessorImplPtr> AccStorage,
                 std::vector<std::shared_ptr<const void>> SharedPtrStorage,
                 std::vector<Requirement *> Requirements,
-                std::vector<detail::EventImplPtr> Events, CGTYPE Type)
+                std::vector<detail::EventImplPtr> Events, CGTYPE Type,
+                detail::code_location loc = {})
       : CG(Type, std::move(ArgsStorage), std::move(AccStorage),
            std::move(SharedPtrStorage), std::move(Requirements),
-           std::move(Events)),
+           std::move(Events), std::move(loc)),
         MInteropTask(std::move(InteropTask)) {}
 };
 

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

@@ -16,24 +16,84 @@
 #include <CL/cl.h>
 #include <CL/cl_ext.h>
 #include <CL/cl_ext_intel.h>
+
+#include <cstdint>
 #include <string>
 #include <type_traits>
 
 #define STRINGIFY_LINE_HELP(s) #s
 #define STRINGIFY_LINE(s) STRINGIFY_LINE_HELP(s)
 
+// Default signature enables the passing of user code location information to
+// public methods as a default argument. If the end-user wants to disable the
+// code location information, they must compile the code with
+// -DDISABLE_SYCL_INSTRUMENTATION_METADATA flag
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+// We define a sycl stream name and this will
+// be used by the instrumentation framework
+constexpr const char *SYCL_STREAM_NAME = "sycl";
+// Data structure that captures the user code
+// location information using the builtin capabilities
+// of the compiler
+struct code_location {
+#ifdef _MSC_VER
+  // Since MSVC does not support the required builtins, we
+  // implement the version with "unknown"s which is handled
+  // correctly by the instrumentation
+  static constexpr code_location current(const char *fileName = nullptr,
+                                         const char *funcName = nullptr,
+                                         unsigned long lineNo = 0,
+                                         unsigned long columnNo = 0) noexcept {
+    return code_location(fileName, funcName, lineNo, columnNo);
+  }
+#else
+  static constexpr code_location
+  current(const char *fileName = __builtin_FILE(),
+          const char *funcName = __builtin_FUNCTION(),
+          unsigned long lineNo = __builtin_LINE(),
+          unsigned long columnNo = 0) noexcept {
+    return code_location(fileName, funcName, lineNo, columnNo);
+  }
+#endif
+
+  constexpr code_location(const char *file, const char *func, int line,
+                          int col) noexcept
+      : MFileName(file), MFunctionName(func), MLineNo(line), MColumnNo(col) {}
+
+  constexpr code_location() noexcept
+      : MFileName(nullptr), MFunctionName(nullptr), MLineNo(0), MColumnNo(0) {}
+
+  constexpr unsigned long lineNumber() const noexcept { return MLineNo; }
+  constexpr unsigned long columnNumber() const noexcept { return MColumnNo; }
+  constexpr const char *fileName() const noexcept { return MFileName; }
+  constexpr const char *functionName() const noexcept { return MFunctionName; }
+
+private:
+  const char *MFileName;
+  const char *MFunctionName;
+  unsigned long MLineNo;
+  unsigned long MColumnNo;
+};
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
 const char *stringifyErrorCode(cl_int error);
 
-static inline std::string codeToString(cl_int code){
-  return std::string(std::to_string(code) + " (" +
-         stringifyErrorCode(code) + ")");
+static inline std::string codeToString(cl_int code) {
+  return std::string(std::to_string(code) + " (" + stringifyErrorCode(code) +
+                     ")");
 }
 
-}}} // __SYCL_INLINE_NAMESPACE(cl)::sycl::detail
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
 
 #ifdef __SYCL_DEVICE_ONLY__
 // TODO remove this when 'assert' is supported in device code

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

@@ -18,6 +18,13 @@
 #include <cassert>
 #include <string>
 
+#ifdef XPTI_ENABLE_INSTRUMENTATION
+// Forward declarations
+namespace xpti {
+struct trace_event_data_t;
+}
+#endif
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {

sycl/include/CL/sycl/handler.hpp

@@ -195,8 +195,9 @@ class handler {
   /// It's expected that the method is the latest method executed before
   /// object destruction.
   ///
-  /// \return a SYCL event object representing the command group.
-  event finalize();
+  /// \param Payload contains the code location of user code
+  /// \return a SYCL event object representing the command group
+  event finalize(const cl::sycl::detail::code_location &Payload = {});
 
   /// Saves streams associated with this handler.
   ///