diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index 129eba7543..11df6ee61c 100644
--- a/.github/CODEOWNERS
+++ b/.github/CODEOWNERS
@@ -77,3 +77,6 @@ SYCL/BFloat16 @AlexeySotkin @MrSidims
 
 # Deprecated features
 SYCL/DeprecatedFeatures @intel/llvm-reviewers-runtime
+
+# XPTI and XPTI Framework
+SYCL/XPTI @intel/llvm-reviewers-runtime
diff --git a/SYCL/XPTI/Inputs/test_collector.cpp b/SYCL/XPTI/Inputs/test_collector.cpp
new file mode 100644
index 0000000000..bd6f7fd70d
--- /dev/null
+++ b/SYCL/XPTI/Inputs/test_collector.cpp
@@ -0,0 +1,138 @@
+#include "xpti/xpti_trace_framework.hpp"
+
+#include <iostream>
+#include <mutex>
+#include <string_view>
+
+std::mutex GMutex;
+
+XPTI_CALLBACK_API void syclCallback(uint16_t, xpti::trace_event_data_t *,
+                                    xpti::trace_event_data_t *, uint64_t,
+                                    const void *);
+XPTI_CALLBACK_API void syclPiCallback(uint16_t, xpti::trace_event_data_t *,
+                                      xpti::trace_event_data_t *, uint64_t,
+                                      const void *);
+
+XPTI_CALLBACK_API void xptiTraceInit(unsigned int MajorVersion,
+                                     unsigned int MinorVersion,
+                                     const char *VersionStr,
+                                     const char *StreamName) {
+  std::cout << "xptiTraceInit: Stream Name = " << StreamName << "\n";
+  std::string_view NameView{StreamName};
+
+  if (NameView == "sycl.pi") {
+    uint8_t StreamID = xptiRegisterStream(StreamName);
+    xptiRegisterCallback(
+        StreamID,
+        static_cast<uint16_t>(xpti::trace_point_type_t::function_begin),
+        syclPiCallback);
+    xptiRegisterCallback(
+        StreamID,
+        static_cast<uint16_t>(xpti::trace_point_type_t::function_with_args_end),
+        syclPiCallback);
+  }
+  if (NameView == "sycl") {
+    uint8_t StreamID = xptiRegisterStream(StreamName);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::graph_create),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::node_create),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::edge_create),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::task_begin),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::task_end),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::signal),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID,
+        static_cast<uint16_t>(xpti::trace_point_type_t::barrier_begin),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::barrier_end),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::wait_begin),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::wait_end),
+        syclCallback);
+    xptiRegisterCallback(
+        StreamID, static_cast<uint16_t>(xpti::trace_point_type_t::signal),
+        syclCallback);
+  }
+}
+
+XPTI_CALLBACK_API void xptiTraceFinish(const char *streamName) {
+  std::cout << "xptiTraceFinish: Stream Name = " << streamName << "\n";
+}
+
+XPTI_CALLBACK_API void syclPiCallback(uint16_t TraceType,
+                                      xpti::trace_event_data_t *,
+                                      xpti::trace_event_data_t *, uint64_t,
+                                      const void *UserData) {
+  std::lock_guard Lock{GMutex};
+  auto Type = static_cast<xpti::trace_point_type_t>(TraceType);
+  const char *funcName = static_cast<const char *>(UserData);
+  if (Type == xpti::trace_point_type_t::function_begin) {
+    std::cout << "PI Call Begin : ";
+  } else if (Type == xpti::trace_point_type_t::function_end) {
+    std::cout << "PI Call End : ";
+  }
+  std::cout << funcName << "\n";
+}
+
+XPTI_CALLBACK_API void syclCallback(uint16_t TraceType,
+                                    xpti::trace_event_data_t *,
+                                    xpti::trace_event_data_t *Event, uint64_t,
+                                    const void *UserData) {
+  std::lock_guard Lock{GMutex};
+  auto Type = static_cast<xpti::trace_point_type_t>(TraceType);
+  switch (Type) {
+  case xpti::trace_point_type_t::graph_create:
+    std::cout << "Graph create\n";
+    break;
+  case xpti::trace_point_type_t::node_create:
+    std::cout << "Node create\n";
+    break;
+  case xpti::trace_point_type_t::edge_create:
+    std::cout << "Edge create\n";
+    break;
+  case xpti::trace_point_type_t::task_begin:
+    std::cout << "Task begin\n";
+    break;
+  case xpti::trace_point_type_t::task_end:
+    std::cout << "Task end\n";
+    break;
+  case xpti::trace_point_type_t::signal:
+    std::cout << "Signal\n";
+    break;
+  case xpti::trace_point_type_t::wait_begin:
+    std::cout << "Wait begin\n";
+    break;
+  case xpti::trace_point_type_t::wait_end:
+    std::cout << "Wait end\n";
+    break;
+  case xpti::trace_point_type_t::barrier_begin:
+    std::cout << "Barrier begin\n";
+    break;
+  case xpti::trace_point_type_t::barrier_end:
+    std::cout << "Barrier end\n";
+    break;
+  default:
+    std::cout << "Unknown tracepoint\n";
+  }
+
+  xpti::metadata_t *Metadata = xptiQueryMetadata(Event);
+  for (auto &Item : *Metadata) {
+    std::cout << "  " << xptiLookupString(Item.first) << " : "
+              << xptiLookupString(Item.second) << "\n";
+  }
+}
diff --git a/SYCL/XPTI/basic_event_collection.cpp b/SYCL/XPTI/basic_event_collection.cpp
new file mode 100644
index 0000000000..9e8907bcb3
--- /dev/null
+++ b/SYCL/XPTI/basic_event_collection.cpp
@@ -0,0 +1,121 @@
+// REQUIRES: xptifw, opencl
+// RUN: %clangxx %s -DXPTI_COLLECTOR -DXPTI_CALLBACK_API_EXPORTS %xptifw_lib %shared_lib %fPIC %cxx_std_optionc++17 -o %t_collector.dll
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env XPTI_TRACE_ENABLE=1 env XPTI_FRAMEWORK_DISPATCHER=%xptifw_dispatcher env XPTI_SUBSCRIBERS=%t_collector.dll env SYCL_DEVICE_FILTER=opencl %t.out | FileCheck %s 2>&1
+
+#ifdef XPTI_COLLECTOR
+
+#include "Inputs/test_collector.cpp"
+
+#else
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q{sycl::default_selector{}};
+
+  auto Ptr = sycl::malloc_device<int>(1, Q);
+
+  auto Evt1 = Q.single_task([=]() { Ptr[0] = 1; });
+
+  auto Evt2 = Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(Evt1);
+    CGH.single_task([=]() { Ptr[0]++; });
+  });
+
+  Evt2.wait();
+
+  int Res = 0;
+  Q.memcpy(&Res, Ptr, 1);
+  Q.wait();
+
+  assert(Res == 2);
+
+  return 0;
+}
+
+#endif
+
+// CHECK: xptiTraceInit: Stream Name = sycl
+// CHECK-NEXT: Graph create
+// CHECK-NEXT: xptiTraceInit: Stream Name = sycl.pi
+// CHECK-NEXT: xptiTraceInit: Stream Name = sycl.pi.debug
+// CHECK-NEXT: PI Call Begin : piPlatformsGet
+// CHECK-NEXT: PI Call Begin : piPlatformsGet
+// CHECK-NEXT: PI Call Begin : piDevicesGet
+// CHECK-NEXT: PI Call Begin : piDevicesGet
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK-NEXT: PI Call Begin : piDeviceRetain
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
+// CHECK-NEXT: PI Call Begin : piPlatformGetInfo
+// CHECK-NEXT: PI Call Begin : piDeviceRelease
+// CHECK:      PI Call Begin : piContextCreate
+// CHECK-NEXT: PI Call Begin : piQueueCreate
+// CHECK-NEXT: PI Call Begin : piextUSMDeviceAlloc
+// CHECK-NEXT: PI Call Begin : piextDeviceSelectBinary
+// CHECK-NEXT: PI Call Begin : piDeviceGetInfo
+// CHECK:      PI Call Begin : piKernelCreate
+// CHECK-NEXT: PI Call Begin : piKernelSetExecInfo
+// CHECK-NEXT: PI Call Begin : piextKernelSetArgPointer
+// CHECK-NEXT: PI Call Begin : piKernelGetGroupInfo
+// CHECK-NEXT: PI Call Begin : piEnqueueKernelLaunch
+// CHECK-NEXT: Node create
+// CHECK-NEXT:   sym_line_no : 21
+// CHECK-NEXT:   sym_source_file_name : {{.*}}
+// CHECK-NEXT:   sym_function_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   from_source : false
+// CHECK-NEXT:   kernel_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: Node create
+// CHECK-NEXT:   kernel_name : virtual_node[{{.*}}]
+// CHECK-NEXT: Edge create
+// CHECK-NEXT:   event : Event[{{.*}}]
+// CHECK-NEXT: Task begin
+// CHECK-NEXT:   sym_line_no : 21
+// CHECK-NEXT:   sym_source_file_name : {{.*}}
+// CHECK-NEXT:   sym_function_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   from_source : false
+// CHECK-NEXT:   kernel_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: PI Call Begin : piKernelCreate
+// CHECK-NEXT: PI Call Begin : piKernelSetExecInfo
+// CHECK-NEXT: PI Call Begin : piextKernelSetArgPointer
+// CHECK-NEXT: PI Call Begin : piKernelGetGroupInfo
+// CHECK-NEXT: PI Call Begin : piEnqueueKernelLaunch
+// CHECK-NEXT: Signal
+// CHECK-NEXT:   sym_line_no : 21
+// CHECK-NEXT:   sym_source_file_name : {{.*}}
+// CHECK-NEXT:   sym_function_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   from_source : false
+// CHECK-NEXT:   kernel_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: Task end
+// CHECK-NEXT:   sym_line_no : 21
+// CHECK-NEXT:   sym_source_file_name : {{.*}}
+// CHECK-NEXT:   sym_function_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   from_source : false
+// CHECK-NEXT:   kernel_name : typeinfo name for main::{lambda(cl::sycl::handler&)#1}::operator()(cl::sycl::handler&) const::{lambda()#1}
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: Wait begin
+// CHECK-NEXT: PI Call Begin : piEventsWait
+// CHECK-NEXT: Wait end
+// CHECK-NEXT: PI Call Begin : piextUSMEnqueueMemcpy
+// CHECK-NEXT: PI Call Begin : piEventRelease
+// CHECK-NEXT: Wait begin
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: PI Call Begin : piQueueFinish
+// CHECK-NEXT: Wait end
+// CHECK-NEXT:   sycl_device : {{.*}}
+// CHECK-NEXT: PI Call Begin : piEventRelease
+// CHECK-NEXT: PI Call Begin : piEventRelease
+// CHECK-NEXT: PI Call Begin : piQueueRelease
+// CHECK-NEXT: PI Call Begin : piContextRelease
+// CHECK-NEXT: PI Call Begin : piKernelRelease
+// CHECK-NEXT: PI Call Begin : piKernelRelease
+// CHECK-NEXT: PI Call Begin : piProgramRelease
+// CHECK-NEXT: PI Call Begin : piDeviceRelease
+// CHECK-NEXT: PI Call Begin : piTearDown
diff --git a/SYCL/lit.cfg.py b/SYCL/lit.cfg.py
index 8acb369117..90e07579ca 100644
--- a/SYCL/lit.cfg.py
+++ b/SYCL/lit.cfg.py
@@ -142,12 +142,16 @@
     config.substitutions.append( ('%include_option',  '/FI' ) )
     config.substitutions.append( ('%debug_option',  '/DEBUG' ) )
     config.substitutions.append( ('%cxx_std_option',  '/std:' ) )
+    config.substitutions.append( ('%fPIC', '') )
+    config.substitutions.append( ('%shared_lib', '/LD') )
 else:
     config.substitutions.append( ('%sycl_options', ' -lsycl -I' +
                                 config.sycl_include + ' -I' + os.path.join(config.sycl_include, 'sycl')) )
     config.substitutions.append( ('%include_option',  '-include' ) )
     config.substitutions.append( ('%debug_option',  '-g' ) )
     config.substitutions.append( ('%cxx_std_option',  '-std=' ) )
+    config.substitutions.append( ('%fPIC', '-fPIC') )
+    config.substitutions.append( ('%shared_lib', '-shared') )
 
 if not config.gpu_aot_target_opts:
     config.gpu_aot_target_opts = '"-device *"'
@@ -326,6 +330,34 @@
 if find_executable('sycl-ls'):
     config.available_features.add('sycl-ls')
 
+# TODO properly set XPTIFW include and runtime dirs
+xptifw_lib_dir = os.path.join(config.dpcpp_root_dir, 'lib')
+xptifw_dispatcher = ""
+if platform.system() == "Linux":
+    xptifw_dispatcher = os.path.join(xptifw_lib_dir, 'libxptifw.so')
+elif platform.system() == "Windows":
+    xptifw_dispatcher = os.path.join(config.dpcpp_root_dir, 'bin', 'xptifw.dll')
+xptifw_includes = os.path.join(config.dpcpp_root_dir, 'include')
+if os.path.exists(xptifw_lib) and os.path.exists(os.path.join(xptifw_includes, 'xpti', 'xpti_trace_framework.h')):
+    config.available_features.add('xptifw')
+    config.substitutions.append(('%xptifw_dispatcher', xptifw_dispatcher))
+    if platform.system() == "Linux":
+        config.substitutions.append(('%xptifw_lib', " {}/xptifw.lib".format(xptifw_lib_dir)))
+    elif platform.system() == "Windows":
+        config.substitutions.append(('%xptifw_lib', "-L{} -I{} -lxptifw".format(xptifw_lib_dir, xptifw_includes)))
+
+
+llvm_tools = ["llvm-spirv", "llvm-link"]
+for llvm_tool in llvm_tools:
+  llvm_tool_path = find_executable(llvm_tool)
+  if llvm_tool_path:
+    lit_config.note("Found " + llvm_tool)
+    config.available_features.add(llvm_tool)
+    config.substitutions.append( ('%' + llvm_tool.replace('-', '_'),
+                                  os.path.realpath(llvm_tool_path)) )
+  else:
+    lit_config.warning("Can't find " + llvm_tool)
+
 if find_executable('cmc'):
     config.available_features.add('cm-compiler')