Merge from 'sycl' to 'sycl-web'

Author: iclsrc

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -17,6 +17,7 @@
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/Path.h"
+#include <algorithm>
 #include <sstream>
 
 using namespace clang::driver;
@@ -311,6 +312,21 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
       if (SanitizeVal == "address")
         addLibraries(SYCLDeviceSanitizerLibs);
     }
+  } else {
+    // User can pass -fsanitize=address to device compiler via
+    // -Xsycl-target-frontend, sanitize device library must be
+    // linked with user's device image if so.
+    bool IsDeviceAsanEnabled = false;
+    auto SyclFEArg = Args.getAllArgValues(options::OPT_Xsycl_frontend);
+    IsDeviceAsanEnabled = (std::count(SyclFEArg.begin(), SyclFEArg.end(),
+                                      "-fsanitize=address") > 0);
+    if (!IsDeviceAsanEnabled) {
+      auto SyclFEArgEq = Args.getAllArgValues(options::OPT_Xsycl_frontend_EQ);
+      IsDeviceAsanEnabled = (std::count(SyclFEArgEq.begin(), SyclFEArgEq.end(),
+                                        "-fsanitize=address") > 0);
+    }
+    if (IsDeviceAsanEnabled)
+      addLibraries(SYCLDeviceSanitizerLibs);
   }
 #endif
   return LibraryList;

clang/test/Driver/sycl-device-lib.cpp

@@ -216,6 +216,10 @@
 /// test behavior of libsycl-sanitizer.o linking when -fsanitize=address is available
 // RUN: %clangxx -fsycl %s --sysroot=%S/Inputs/SYCL -fsanitize=address -### 2>&1 \
 // RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN: %clangxx -fsycl %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend -fsanitize=address -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN: %clangxx -fsycl %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend=spir64 -fsanitize=address -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
 // SYCL_DEVICE_LIB_SANITIZER: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64-unknown-unknown" "-input={{.*}}libsycl-crt.o" "-output={{.*}}libsycl-crt-{{.*}}.o" "-unbundle"
 // SYCL_DEVICE_LIB_SANITIZER-NEXT: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64-unknown-unknown" "-input={{.*}}libsycl-complex.o" "-output={{.*}}libsycl-complex-{{.*}}.o" "-unbundle"
 // SYCL_DEVICE_LIB_SANITIZER-NEXT: clang-offload-bundler{{.*}} "-type=o" "-targets=sycl-spir64-unknown-unknown" "-input={{.*}}libsycl-complex-fp64.o" "-output={{.*}}libsycl-complex-fp64-{{.*}}.o" "-unbundle"

llvm/test/CodeGen/Generic/fp-builtin-intrinsics-svml.ll

@@ -241,11 +241,22 @@ created on UR command-buffer enqueue.
 
 There is also a *WaitEvent* used by the `ur_exp_command_buffer_handle_t` class
 in the prefix to wait on any dependencies passed in the enqueue wait-list.
-This WaitEvent is reset at the end of the suffix, along with reset commands
-to reset the L0 events used to implement the UR sync-points back to the
-non-signaled state.
-
-![L0 command-buffer diagram](images/L0_UR_command-buffer.svg)
+This WaitEvent is reset in the suffix.
+
+A command-buffer is expected to be submitted multiple times. Consequently,
+we need to ensure that L0 events associated with graph commands have not
+been signaled by a previous execution. These events are therefore reset to the
+non-signaled state before running the actual graph associated commands. Note
+that this reset is performed in the prefix and not in the suffix to avoid
+additional synchronization w.r.t profiling data extraction.
+
+If a command-buffer is about to be submitted to a queue with the profiling
+property enabled, an extra command that copies timestamps of L0 events
+associated with graph commands into a dedicated memory which is attached to the
+returned UR event. This memory stores the profiling information that
+corresponds to the current submission of the command-buffer.
+
+![L0 command-buffer diagram](images/L0_UR_command-buffer-v3.jpg)
 
 For a call to `urCommandBufferEnqueueExp` with an `event_list` *EL*,
 command-buffer *CB*, and return event *RE* our implementation has to submit two

sycl/doc/design/CommandGraph.md

@@ -52,14 +52,12 @@ This extension also depends on the following other SYCL extensions:
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.
-Interfaces defined in this specification may not be implemented yet or may be
-in a preliminary state.
-The specification itself may also change in incompatible ways before it is
-finalized.
-*Shipping software products should not rely on APIs defined in this
-specification.*
+This is an experimental extension specification, intended to provide early 
+access to features and gather community feedback. Interfaces defined in 
+this specification are implemented in DPC++, but they are not finalized 
+and may change incompatibly in future versions of DPC++ without prior notice. 
+*Shipping software products should not rely on APIs defined in 
+this specification.*
 
 
 == Overview

sycl/doc/design/images/L0_UR_command-buffer-v3.jpg

@@ -25,6 +25,7 @@
 #include <sycl/ext/oneapi/properties/property_value.hpp> // and log
 
 #include <array>       // for array
+#include <cstddef>     // for std::byte
 #include <cstring>     // for size_t, memcpy
 #include <functional>  // for function
 #include <iterator>    // for distance
@@ -46,7 +47,7 @@ auto get_native(const kernel_bundle<State> &Obj)
 namespace detail {
 class kernel_id_impl;
 class kernel_impl;
-}
+} // namespace detail
 
 template <typename KernelName> kernel_id get_kernel_id();
 
@@ -886,11 +887,18 @@ __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE,
 /////////////////////////
 // syclex::create_kernel_bundle_from_source
 /////////////////////////
+
+__SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::string &Source);
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
 __SYCL_EXPORT kernel_bundle<bundle_state::ext_oneapi_source>
-create_kernel_bundle_from_source(
-    const context &SyclContext,
-    sycl::ext::oneapi::experimental::source_language Language,
-    const std::string &Source);
+create_kernel_bundle_from_source(const context &SyclContext,
+                                 source_language Language,
+                                 const std::vector<std::byte> &Bytes);
+#endif
 
 /////////////////////////
 // syclex::build(source_kb) => exe_kb

sycl/doc/design/images/L0_UR_command-buffer.svg

@@ -20,7 +20,7 @@ enum class bundle_state : char {
 
 namespace ext::oneapi::experimental {
 
-enum class source_language : int { opencl = 0 /* sycl , spir-v, cuda */ };
+enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ };
 
 } // namespace ext::oneapi::experimental
 

sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc renamed to sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc

@@ -794,11 +794,7 @@ inline pi_result piTearDown(void *PluginParameter) {
   return PI_SUCCESS;
 }
 
-///////////////////////////////////////////////////////////////////////////////
-// Platform
-inline pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms,
-                                pi_uint32 *NumPlatforms) {
-
+inline pi_result PiGetAdapter(ur_adapter_handle_t &adapter) {
   // We're not going through the UR loader so we're guaranteed to have exactly
   // one adapter (whichever is statically linked). The PI plugin for UR has its
   // own implementation of piPlatformsGet.
@@ -809,9 +805,23 @@ inline pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms,
                  [&Ret]() { Ret = urAdapterGet(1, &Adapter, nullptr); });
   HANDLE_ERRORS(Ret);
 
+  adapter = Adapter;
+
+  return PI_SUCCESS;
+}
+
+///////////////////////////////////////////////////////////////////////////////
+// Platform
+inline pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms,
+                                pi_uint32 *NumPlatforms) {
+  ur_adapter_handle_t adapter = nullptr;
+  if (auto res = PiGetAdapter(adapter); res != PI_SUCCESS) {
+    return res;
+  }
+
   auto phPlatforms = reinterpret_cast<ur_platform_handle_t *>(Platforms);
   HANDLE_ERRORS(
-      urPlatformGet(&Adapter, 1, NumEntries, phPlatforms, NumPlatforms));
+      urPlatformGet(&adapter, 1, NumEntries, phPlatforms, NumPlatforms));
   return PI_SUCCESS;
 }
 
@@ -838,6 +848,12 @@ piextPlatformCreateWithNativeHandle(pi_native_handle NativeHandle,
   PI_ASSERT(Platform, PI_ERROR_INVALID_PLATFORM);
   PI_ASSERT(NativeHandle, PI_ERROR_INVALID_VALUE);
 
+  ur_adapter_handle_t adapter = nullptr;
+  if (auto res = PiGetAdapter(adapter); res != PI_SUCCESS) {
+    return res;
+  }
+  (void)adapter;
+
   ur_platform_handle_t UrPlatform{};
   ur_native_handle_t UrNativeHandle =
       reinterpret_cast<ur_native_handle_t>(NativeHandle);

sycl/include/sycl/kernel_bundle.hpp

@@ -278,17 +278,33 @@ void event_impl::checkProfilingPreconditions() const {
         "Profiling information is unavailable as the queue associated with "
         "the event does not have the 'enable_profiling' property.");
   }
-  if (MEventFromSubmitedExecCommandBuffer) {
-    throw sycl::exception(make_error_code(sycl::errc::invalid),
-                          "Profiling information is unavailable for events "
-                          "returned by a graph submission.");
-  }
 }
 
 template <>
 uint64_t
 event_impl::get_profiling_info<info::event_profiling::command_submit>() {
   checkProfilingPreconditions();
+  // The delay between the submission and the actual start of a CommandBuffer
+  // can be short. Consequently, the submission time, which is based on
+  // an estimated clock and not on the real device clock, may be ahead of the
+  // start time, which is based on the actual device clock.
+  // MSubmitTime is set in a critical performance path.
+  // Force reading the device clock when setting MSubmitTime may deteriorate
+  // the performance.
+  // Since submit time is an estimated time, we implement this little hack
+  // that allows all profiled time to be meaningful.
+  // (Note that the observed time deviation between the estimated clock and
+  // the real device clock is typically less than 0.5ms. The approximation we
+  // made by forcing the re-sync of submit time to start time is less than
+  // 0.5ms. These timing values were obtained empirically using an integrated
+  // Intel GPU).
+  if (MEventFromSubmittedExecCommandBuffer && !MHostEvent && MEvent) {
+    uint64_t StartTime =
+        get_event_profiling_info<info::event_profiling::command_start>(
+            this->getHandleRef(), this->getPlugin());
+    if (StartTime < MSubmitTime)
+      MSubmitTime = StartTime;
+  }
   return MSubmitTime;
 }
 

sycl/include/sycl/kernel_bundle_enums.hpp

@@ -282,12 +282,12 @@ class event_impl {
     return MGraph.lock();
   }
 
-  void setEventFromSubmitedExecCommandBuffer(bool value) {
-    MEventFromSubmitedExecCommandBuffer = value;
+  void setEventFromSubmittedExecCommandBuffer(bool value) {
+    MEventFromSubmittedExecCommandBuffer = value;
   }
 
-  bool isEventFromSubmitedExecCommandBuffer() const {
-    return MEventFromSubmitedExecCommandBuffer;
+  bool isEventFromSubmittedExecCommandBuffer() const {
+    return MEventFromSubmittedExecCommandBuffer;
   }
 
 protected:
@@ -340,8 +340,8 @@ class event_impl {
   /// Store the command graph associated with this event, if any.
   /// This event is also be stored in the graph so a weak_ptr is used.
   std::weak_ptr<ext::oneapi::experimental::detail::graph_impl> MGraph;
-  /// Indicates that the event results from a command graph submission
-  bool MEventFromSubmitedExecCommandBuffer = false;
+  /// Indicates that the event results from a command graph submission.
+  bool MEventFromSubmittedExecCommandBuffer = false;
 
   // If this event represents a submission to a
   // sycl::detail::pi::PiExtCommandBuffer the sync point for that submission is

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -757,7 +757,6 @@ exec_graph_impl::enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue,
     auto NewEvent = std::make_shared<sycl::detail::event_impl>(Queue);
     NewEvent->setContextImpl(Queue->getContextImplPtr());
     NewEvent->setStateIncomplete();
-    NewEvent->setEventFromSubmitedExecCommandBuffer(true);
     return NewEvent;
   });
 
@@ -840,7 +839,7 @@ exec_graph_impl::enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue,
         NewEvent = sycl::detail::Scheduler::getInstance().addCG(
             std::move(CommandGroup), Queue);
       }
-
+      NewEvent->setEventFromSubmittedExecCommandBuffer(true);
     } else if ((CurrentPartition->MSchedule.size() > 0) &&
                (CurrentPartition->MSchedule.front()->MCGType ==
                 sycl::detail::CG::CGTYPE::CodeplayHostTask)) {