Merge from 'sycl' to 'sycl-web'

Author: iclsrc

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.cpp

@@ -126,6 +126,7 @@ void SPIRVDecorate::encode(spv_ostream &O) const {
     break;
   case internal::DecorationFuncParamDescINTEL:
     SPIRVDecorateFuncParamDescAttr::encodeLiterals(Encoder, Literals);
+    break;
   case internal::DecorationHostAccessINTEL:
     SPIRVDecorateHostAccessINTELLegacy::encodeLiterals(Encoder, Literals);
     break;
@@ -163,6 +164,7 @@ void SPIRVDecorate::decode(std::istream &I) {
     break;
   case internal::DecorationFuncParamDescINTEL:
     SPIRVDecorateFuncParamDescAttr::decodeLiterals(Decoder, Literals);
+    break;
   case internal::DecorationHostAccessINTEL:
     SPIRVDecorateHostAccessINTELLegacy::decodeLiterals(Decoder, Literals);
     break;

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1280,7 +1280,8 @@ struct FunctionStackPoisoner : public InstVisitor<FunctionStackPoisoner> {
 static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM) {
   SmallVector<Function *> SpirFixupFuncs;
   for (Function &F : M) {
-    if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
+    if (F.getCallingConv() == CallingConv::SPIR_KERNEL &&
+        F.hasFnAttribute(Attribute::SanitizeAddress)) {
       SpirFixupFuncs.emplace_back(&F);
     }
   }
@@ -1296,9 +1297,7 @@ static void ExtendSpirKernelArgs(Module &M, FunctionAnalysisManager &FAM) {
       Types.push_back(I->getType());
     }
 
-    // New argument type: uintptr_t as(1)*, as it's allocated in USM buffer, and
-    // it can also be treated as a pointer point to the base address of private
-    // shadow memory
+    // New argument: uintptr_t as(1)*, which is allocated in shared USM buffer
     Types.push_back(IntptrTy->getPointerTo(kSpirOffloadGlobalAS));
 
     FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false);
@@ -1413,6 +1412,21 @@ PreservedAnalyses AddressSanitizerPass::run(Module &M,
       ClUseStackSafety ? &MAM.getResult<StackSafetyGlobalAnalysis>(M) : nullptr;
 
   if (Triple(M.getTargetTriple()).isSPIR()) {
+    bool HasESIMDKernel = false;
+
+    // ESIMD kernel doesn't support noinline functions, so we can't
+    // support sanitizer for it
+    for (Function &F : M)
+      if (F.hasMetadata("sycl_explicit_simd")) {
+        F.removeFnAttr(Attribute::SanitizeAddress);
+        HasESIMDKernel = true;
+      }
+
+    // FIXME: we can't check if the kernel is ESIMD kernel at UR, so we
+    // have to disable ASan completely in this case
+    if (HasESIMDKernel)
+      return PreservedAnalyses::all();
+
     ExtendSpirKernelArgs(M, FAM);
   }
 

llvm/test/Instrumentation/AddressSanitizer/SPIRV/sycl_esimd.ll

@@ -0,0 +1,17 @@
+; RUN: opt < %s -passes=asan -asan-instrumentation-with-call-threshold=0 -asan-stack=0 -asan-globals=0 -S | FileCheck %s
+
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+@__spirv_BuiltInGlobalInvocationId = external addrspace(1) constant <3 x i64>
+
+; Function Attrs: sanitize_address
+define spir_kernel void @esimd_kernel() #0 !sycl_explicit_simd !1 {
+entry:
+  %0 = load i64, ptr addrspace(1) getelementptr inbounds (i8, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, i64 8), align 8
+  ret void
+}
+; CHECK-NOT: {{ sanitize_address }}
+
+attributes #0 = { sanitize_address }
+!1 = !{}

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -1736,15 +1736,23 @@ passed an invalid event.
 The new handler methods, and queue shortcuts, defined by
 link:../supported/sycl_ext_oneapi_enqueue_barrier.asciidoc[sycl_ext_oneapi_enqueue_barrier]
 can only be used in graph nodes created using the Record & Replay API, as
-barriers rely on events to enforce dependencies. For barriers with an empty
-wait list parameter, the semantics are that the barrier node being added to
-will depend on all the existing graph leaf nodes, not only the leaf nodes
-that were added from the queue being recorded.
+barriers rely on events to enforce dependencies.
 
 A synchronous exception will be thrown with error code `invalid` if a user
-tries to add them to a graph using the Explicit API. Empty nodes created with
-the `node::depends_on_all_leaves` property can be used instead of barriers when
-a user is building a graph with the explicit API.
+tries to add a barrier command to a graph using the explicit API. Empty nodes
+created with the `node::depends_on_all_leaves` property can be used instead of
+barriers when a user is building a graph with the explicit API.
+
+The semantics of barriers are defined in `sycl_ext_oneapi_enqueue_barrier` for
+a single command-queue, and correlate as follows to a graph that may contain
+nodes that are recorded from multiple queues and/or added by the explicit API:
+
+* Barriers with an empty wait list parameter will only depend on the leaf nodes
+  that were added to the graph from the queue the barrier command is being
+  recorded from.
+
+* The only commands which have an implicit dependency on the barrier command
+  are those recorded from the same queue the barrier command was submitted to.
 
 ==== sycl_ext_oneapi_memcpy2d
 

sycl/doc/syclcompat/README.md

@@ -78,18 +78,24 @@ those to learn to use the library.
 
 SYCLcompat provides a `dim3` class akin to that of CUDA or HIP programming
 models. `dim3` encapsulates other languages iteration spaces that are
-represented with coordinate letters (x, y, z).
+represented with coordinate letters (x, y, z). In SYCL, the fastest-moving
+dimension is the one with the highest index, e.g. in a SYCL 2D range iteration
+space, there are two dimensions, 0 and 1, and 1 will be the one that "moves
+faster". For CUDA/HIP, the convention is reversed: `x` is the dimension which
+moves fastest. `syclcompat::dim3` follows this convention, so that
+`syclcompat::dim3(32, 4)` is equivalent to `sycl::range<2>(4, 32)`, and
+`syclcompat::dim3(32, 4, 2)` is equivalent to `sycl::range<3>(2, 4, 32)`.
 
 ```cpp
 namespace syclcompat {
 
 class dim3 {
 public:
-  const size_t x, y, z;
+  unsigned int x, y, z;
   dim3(const sycl::range<3> &r);
   dim3(const sycl::range<2> &r);
   dim3(const sycl::range<1> &r);
-  constexpr dim3(size_t x, size_t y = 1, size_t z = 1);
+  constexpr dim3(unsigned int x = 1, unsigned int y = 1, unsigned int z = 1);
 
   constexpr size_t size();
 
@@ -106,12 +112,10 @@ inline dim3 operator-(const dim3 &a, const dim3 &b);
 } // syclcompat
 ```
 
-In SYCL, the fastest-moving dimension is the one with the highest index, e.g. in
-a SYCL 2D range iteration space, there are two dimensions, 0 and 1, and 1 will
-be the one that "moves faster". The compatibility headers for SYCL offer a
-number of convenience functions that help the mapping between xyz-based
-coordinates to SYCL iteration spaces in the different scopes available. In
-addition to the global range, the following helper functions are also provided:
+The compatibility headers for SYCL offer a number of convenience functions that
+help the mapping between xyz-based coordinates to SYCL iteration spaces in the
+different scopes available. In addition to the global range, the following
+helper functions are also provided:
 
 ``` c++
 namespace syclcompat {

sycl/include/sycl/handler.hpp

@@ -466,7 +466,7 @@ class __SYCL_EXPORT handler {
   /// \param Queue is a SYCL queue.
   /// \param IsHost indicates if this handler is created for SYCL host device.
   /// TODO: Unused. Remove with ABI break.
-  handler(std::shared_ptr<detail::queue_impl> Queue, bool IsHost);
+  handler(std::shared_ptr<detail::queue_impl> Queue, bool /*Unused*/);
 
   /// Constructs SYCL handler from the associated queue and the submission's
   /// primary and secondary queue.
@@ -476,20 +476,20 @@ class __SYCL_EXPORT handler {
   /// \param PrimaryQueue is the primary SYCL queue of the submission.
   /// \param SecondaryQueue is the secondary SYCL queue of the submission. This
   ///        is null if no secondary queue is associated with the submission.
-  /// \param IsHost indicates if this handler is created for SYCL host device.
   /// TODO: Unused. Remove with ABI break.
   handler(std::shared_ptr<detail::queue_impl> Queue,
           std::shared_ptr<detail::queue_impl> PrimaryQueue,
-          std::shared_ptr<detail::queue_impl> SecondaryQueue, bool IsHost);
+          std::shared_ptr<detail::queue_impl> SecondaryQueue,
+          bool /* Unused */);
 
   /// Constructs SYCL handler from queue.
   ///
   /// \param Queue is a SYCL queue.
   /// \param IsHost indicates if this handler is created for SYCL host device.
   /// \param CallerNeedsEvent indicates if the event resulting from this handler
   ///        is needed by the caller.
-  handler(std::shared_ptr<detail::queue_impl> Queue, bool IsHost,
-          bool CallerNeedsEvent);
+  handler(std::shared_ptr<detail::queue_impl> Queue,
+          bool /* ABI break: remove */, bool CallerNeedsEvent);
 
   /// Constructs SYCL handler from the associated queue and the submission's
   /// primary and secondary queue.
@@ -504,8 +504,8 @@ class __SYCL_EXPORT handler {
   ///        is needed by the caller.
   handler(std::shared_ptr<detail::queue_impl> Queue,
           std::shared_ptr<detail::queue_impl> PrimaryQueue,
-          std::shared_ptr<detail::queue_impl> SecondaryQueue, bool IsHost,
-          bool CallerNeedsEvent);
+          std::shared_ptr<detail::queue_impl> SecondaryQueue,
+          bool /* ABI break: remove */, bool CallerNeedsEvent);
 
   /// Constructs SYCL handler from Graph.
   ///
@@ -644,7 +644,7 @@ class __SYCL_EXPORT handler {
   ~handler() = default;
 
   // TODO: Private and unusued. Remove when ABI break is allowed.
-  bool is_host() { return MIsHost; }
+  bool is_host() { return false; }
 
 #ifdef __SYCL_DEVICE_ONLY__
   // In device compilation accessor isn't inherited from host base classes, so
@@ -923,12 +923,6 @@ class __SYCL_EXPORT handler {
         detail::KernelLambdaHasKernelHandlerArgT<KernelType,
                                                  LambdaArgType>::value;
 
-    if (IsCallableWithKernelHandler && MIsHost) {
-      throw sycl::feature_not_supported(
-          "kernel_handler is not yet supported by host device.",
-          PI_ERROR_INVALID_OPERATION);
-    }
-
     KernelType *KernelPtr =
         ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);
 
@@ -1077,8 +1071,7 @@ class __SYCL_EXPORT handler {
   std::enable_if_t<(DimSrc > 0) && (DimDst > 0), bool>
   copyAccToAccHelper(accessor<TSrc, DimSrc, ModeSrc, TargetSrc, IsPHSrc> Src,
                      accessor<TDst, DimDst, ModeDst, TargetDst, IsPHDst> Dst) {
-    if (!MIsHost &&
-        IsCopyingRectRegionAvailable(Src.get_range(), Dst.get_range()))
+    if (IsCopyingRectRegionAvailable(Src.get_range(), Dst.get_range()))
       return false;
 
     range<1> LinearizedRange(Src.size());
@@ -1100,23 +1093,19 @@ class __SYCL_EXPORT handler {
   ///
   /// \param Src is a source SYCL accessor.
   /// \param Dst is a destination SYCL accessor.
+  // ABI break: to remove whole method
   template <typename TSrc, int DimSrc, access::mode ModeSrc,
             access::target TargetSrc, typename TDst, int DimDst,
             access::mode ModeDst, access::target TargetDst,
             access::placeholder IsPHSrc, access::placeholder IsPHDst>
   std::enable_if_t<DimSrc == 0 || DimDst == 0, bool>
-  copyAccToAccHelper(accessor<TSrc, DimSrc, ModeSrc, TargetSrc, IsPHSrc> Src,
-                     accessor<TDst, DimDst, ModeDst, TargetDst, IsPHDst> Dst) {
-    if (!MIsHost)
-      return false;
-
-    single_task<__copyAcc2Acc<TSrc, DimSrc, ModeSrc, TargetSrc, TDst, DimDst,
-                              ModeDst, TargetDst, IsPHSrc, IsPHDst>>(
-        [=]() { *(Dst.get_pointer()) = *(Src.get_pointer()); });
-    return true;
+  copyAccToAccHelper(accessor<TSrc, DimSrc, ModeSrc, TargetSrc, IsPHSrc>,
+                     accessor<TDst, DimDst, ModeDst, TargetDst, IsPHDst>) {
+    return false;
   }
 
 #ifndef __SYCL_DEVICE_ONLY__
+  // ABI break: to remove whole method
   /// Copies the content of memory object accessed by Src into the memory
   /// pointed by Dst.
   ///
@@ -1136,6 +1125,7 @@ class __SYCL_EXPORT handler {
         });
   }
 
+  // ABI break: to remove whole method
   /// Copies 1 element accessed by 0-dimensional accessor Src into the memory
   /// pointed by Dst.
   ///
@@ -1153,6 +1143,7 @@ class __SYCL_EXPORT handler {
         });
   }
 
+  // ABI break: to remove whole method
   /// Copies the memory pointed by Src into the memory accessed by Dst.
   ///
   /// \param Src is a pointer to source memory.
@@ -1170,6 +1161,7 @@ class __SYCL_EXPORT handler {
         });
   }
 
+  // ABI break: to remove whole method
   /// Copies 1 element pointed by Src to memory accessed by 0-dimensional
   /// accessor Dst.
   ///
@@ -2282,7 +2274,7 @@ class __SYCL_EXPORT handler {
     MNDRDesc.set(range<1>{1});
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     setType(detail::CG::Kernel);
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
+    if (!lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
       MKernelName = getKernelName();
     } else
@@ -2319,7 +2311,7 @@ class __SYCL_EXPORT handler {
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     setType(detail::CG::Kernel);
     setNDRangeUsed(false);
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
+    if (!lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
       MKernelName = getKernelName();
     } else
@@ -2359,7 +2351,7 @@ class __SYCL_EXPORT handler {
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     setType(detail::CG::Kernel);
     setNDRangeUsed(false);
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
+    if (!lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
       MKernelName = getKernelName();
     } else
@@ -2398,7 +2390,7 @@ class __SYCL_EXPORT handler {
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     setType(detail::CG::Kernel);
     setNDRangeUsed(true);
-    if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
+    if (!lambdaAndKernelHaveEqualName<NameT>()) {
       extractArgsAndReqs();
       MKernelName = getKernelName();
     } else
@@ -2725,14 +2717,6 @@ class __SYCL_EXPORT handler {
                   "Invalid accessor target for the copy method.");
     static_assert(isValidModeForSourceAccessor(AccessMode),
                   "Invalid accessor mode for the copy method.");
-#ifndef __SYCL_DEVICE_ONLY__
-    if (MIsHost) {
-      // TODO: Temporary implementation for host. Should be handled by memory
-      // manager.
-      copyAccToPtrHost(Src, Dst);
-      return;
-    }
-#endif
     setType(detail::CG::CopyAccToPtr);
 
     detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Src;
@@ -2769,14 +2753,7 @@ class __SYCL_EXPORT handler {
                   "Invalid accessor mode for the copy method.");
     // TODO: Add static_assert with is_device_copyable when vec is
     // device-copyable.
-#ifndef __SYCL_DEVICE_ONLY__
-    if (MIsHost) {
-      // TODO: Temporary implementation for host. Should be handled by memory
-      // manager.
-      copyPtrToAccHost(Src, Dst);
-      return;
-    }
-#endif
+
     setType(detail::CG::CopyPtrToAcc);
 
     detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Dst;
@@ -2890,8 +2867,6 @@ class __SYCL_EXPORT handler {
   fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder, PropertyListT>
            Dst,
        const T &Pattern) {
-    assert(!MIsHost && "fill() should no longer be callable on a host device.");
-
     if (Dst.is_placeholder())
       checkIfPlaceholderIsBoundToHandler(Dst);
 
@@ -3429,7 +3404,7 @@ class __SYCL_EXPORT handler {
   /// Storage for the CG created when handling graph nodes added explicitly.
   std::unique_ptr<detail::CG> MGraphNodeCG;
 
-  bool MIsHost = false;
+  bool MIsHost = false; // ABI break: to remove
 
   detail::code_location MCodeLoc = {};
   bool MIsFinalized = false;

sycl/include/syclcompat/dims.hpp

@@ -30,15 +30,16 @@ namespace syclcompat {
 
 class dim3 {
 public:
-  const size_t x, y, z;
+  unsigned int x, y, z;
 
   dim3(const sycl::range<3> &r) : x(r[2]), y(r[1]), z(r[0]) {}
 
   dim3(const sycl::range<2> &r) : x(r[1]), y(r[0]), z(1) {}
 
   dim3(const sycl::range<1> &r) : x(r[0]), y(1), z(1) {}
 
-  constexpr dim3(size_t x, size_t y = 1, size_t z = 1) : x(x), y(y), z(z) {}
+  constexpr dim3(unsigned int x = 1, unsigned int y = 1, unsigned int z = 1)
+      : x(x), y(y), z(z) {}
 
   constexpr size_t size() const { return x * y * z; }