[SYCL] Optimize TLS usage in NestedCallsTracker and tls_code_loc_t (#18787)

This PR decreases the number of TLS accesses in the `NestedCallsTracker`
and `tls_code_loc_t`. The idea is to cache TLS location in the
reference. As a result, we have only a single lookup for the TLS
location.

Author: vinser52

sycl/include/sycl/detail/common.hpp

@@ -8,6 +8,9 @@
 
 #pragma once
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+#include <sycl/exception.hpp> // for sycl::exception, sycl::errc
+#endif                        // #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
 
@@ -139,6 +142,15 @@ class __SYCL_EXPORT tls_code_loc_t {
   // Used to maintain global state (GCodeLocTLS), so we do not want to copy
   tls_code_loc_t(const tls_code_loc_t &) = delete;
   tls_code_loc_t &operator=(const tls_code_loc_t &) = delete;
+#else
+  tls_code_loc_t &operator=(const tls_code_loc_t &) {
+    // Should never be called. In PREVIEW we marked it as deleted, but
+    // before ABI breaking change we need to keep it for backward compatibility.
+    assert(false && "tls_code_loc_t should not be copied");
+    throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+                          "tls_code_loc_t should not be copied");
+    return *this;
+  }
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
 
   /// If the code location is set up by this instance, reset it.
@@ -152,6 +164,10 @@ class __SYCL_EXPORT tls_code_loc_t {
   bool isToplevel() const { return !MLocalScope; }
 
 private:
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  // Cache the TLS location to decrease amount of TLS accesses.
+  detail::code_location &CodeLocTLSRef;
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
   // The flag that is used to determine if the object is in a local scope or in
   // the top level scope.
   bool MLocalScope = true;

sycl/source/detail/common.cpp

@@ -24,9 +24,18 @@ static thread_local detail::code_location GCodeLocTLS = {};
 /// check and see if code location object is available. If not, continue with
 /// instrumentation as needed
 tls_code_loc_t::tls_code_loc_t()
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    : CodeLocTLSRef(GCodeLocTLS),
+      // Check TLS to see if a previously stashed code_location object is
+      // available; if so, we are in a local scope.
+      MLocalScope(CodeLocTLSRef.fileName() && CodeLocTLSRef.functionName())
+#else
     : // Check TLS to see if a previously stashed code_location object is
       // available; if so, we are in a local scope.
-      MLocalScope(GCodeLocTLS.fileName() && GCodeLocTLS.functionName()) {}
+      MLocalScope(GCodeLocTLS.fileName() && GCodeLocTLS.functionName())
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+{
+}
 
 ur_code_location_t codeLocationCallback(void *) {
   ur_code_location_t codeloc;
@@ -44,24 +53,45 @@ ur_code_location_t codeLocationCallback(void *) {
 /// location has been stashed in the TLS at a higher level. If not, we have the
 /// code location information that must be active for the current calling scope.
 tls_code_loc_t::tls_code_loc_t(const detail::code_location &CodeLoc)
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    : CodeLocTLSRef(GCodeLocTLS),
+      // Check TLS to see if a previously stashed code_location object is
+      // available; if so, then don't overwrite the previous information as we
+      // are still in scope of the instrumented function.
+      MLocalScope(CodeLocTLSRef.fileName() && CodeLocTLSRef.functionName()) {
+  if (!MLocalScope)
+    // Update the TLS information with the code_location information
+    CodeLocTLSRef = CodeLoc;
+#else
     : // Check TLS to see if a previously stashed code_location object is
       // available; if so, then don't overwrite the previous information as we
       // are still in scope of the instrumented function.
       MLocalScope(GCodeLocTLS.fileName() && GCodeLocTLS.functionName()) {
   if (!MLocalScope)
     // Update the TLS information with the code_location information
     GCodeLocTLS = CodeLoc;
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
 }
 
 /// @brief  If we are the top lovel scope,  reset the code location info
 tls_code_loc_t::~tls_code_loc_t() {
   // Only reset the TLS data if the top level function is going out of scope
   if (!MLocalScope) {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+    CodeLocTLSRef = {};
+#else
     GCodeLocTLS = {};
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
   }
 }
 
-const detail::code_location &tls_code_loc_t::query() { return GCodeLocTLS; }
+const detail::code_location &tls_code_loc_t::query() {
+#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
+  return CodeLocTLSRef;
+#else
+  return GCodeLocTLS;
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
+}
 
 } // namespace detail
 } // namespace _V1

sycl/source/detail/queue_impl.cpp

@@ -33,15 +33,19 @@ thread_local bool NestedCallsDetector = false;
 class NestedCallsTracker {
 public:
   NestedCallsTracker() {
-    if (NestedCallsDetector)
+    if (NestedCallsDetectorRef)
       throw sycl::exception(
           make_error_code(errc::invalid),
           "Calls to sycl::queue::submit cannot be nested. Command group "
           "function objects should use the sycl::handler API instead.");
-    NestedCallsDetector = true;
+    NestedCallsDetectorRef = true;
   }
 
-  ~NestedCallsTracker() { NestedCallsDetector = false; }
+  ~NestedCallsTracker() { NestedCallsDetectorRef = false; }
+
+private:
+  // Cache the TLS location to decrease amount of TLS accesses.
+  bool &NestedCallsDetectorRef = NestedCallsDetector;
 };
 
 static std::vector<ur_event_handle_t>

sycl/test/abi/layout_tls_code_loc_t.cpp

@@ -0,0 +1,14 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// REQUIRES: linux
+// UNSUPPORTED: libcxx
+
+// clang-format off
+
+#include <sycl/detail/common.hpp>
+
+void foo(sycl::detail::tls_code_loc_t) {}
+
+// CHECK:      0 | class sycl::detail::tls_code_loc_t
+// CHECK-NEXT: 0 |    _Bool MLocalScope
+// CHECK-NEXT:   | [sizeof=1, dsize=1, align=1,
+// CHECK-NEXT:   |  nvsize=1, nvalign=1]