Merge remote-tracking branch 'origin/sycl' into llvmspirv_pulldown

Author: jsji

.github/CODEOWNERS

@@ -121,6 +121,9 @@ sycl/include/sycl/ext/oneapi/experimental/graph.hpp @intel/sycl-graphs-reviewers
 sycl/source/detail/graph_impl.cpp @intel/sycl-graphs-reviewers
 sycl/source/detail/graph_impl.hpp @intel/sycl-graphs-reviewers
 sycl/unittests/Extensions/CommandGraph.cpp @intel/sycl-graphs-reviewers
+sycl/doc/design/CommandGraph.md @intel/sycl-graphs-reviewers
+sycl/test-e2e/Graph @intel/sycl-graphs-reviewers
+sycl/doc/extensions/**/sycl_ext_oneapi_graph.asciidoc @intel/sycl-graphs-reviewers
 
 # syclcompat library
 sycl/**/syclcompat/ @intel/syclcompat-lib-reviewers

libclc/amdgcn-amdhsa/libspirv/group/collectives.cl

@@ -150,7 +150,7 @@ __CLC_SUBGROUP_COLLECTIVE(FMax, __CLC_MAX, float, f, -INFINITY)
 __CLC_SUBGROUP_COLLECTIVE(FMax, __CLC_MAX, double, d, -INFINITY)
 
 __CLC_SUBGROUP_COLLECTIVE(All, __CLC_AND, bool, a, true)
-__CLC_SUBGROUP_COLLECTIVE(Any, __CLC_OR, bool, a, true)
+__CLC_SUBGROUP_COLLECTIVE(Any, __CLC_OR, bool, a, false)
 
 #undef __CLC_SUBGROUP_COLLECTIVE_BODY
 #undef __CLC_SUBGROUP_COLLECTIVE

test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll renamed to llvm-spirv/test/extensions/KHR/SPV_KHR_non_semantic_info/preserve-all-function-attributes-crash.ll

@@ -14,7 +14,7 @@ deprecated==1.2.13
     # via pygithub
 gitdb==4.0.9
     # via gitpython
-gitpython==3.1.30
+gitpython==3.1.32
     # via -r requirements.txt.in
 idna==3.4
     # via requests

llvm/utils/git/requirements.txt

@@ -1,3 +1,4 @@
+set(SYCL_COMPATH_UNITTEST_GCC_TOOLCHAIN "" CACHE PATH "Path to GCC installation")
 # add_sycl_library_unittest(test_suite_name sycl_extra_flags
 #                           file1.cpp file2.cpp ...)
 #
@@ -26,6 +27,12 @@ macro(add_sycl_library_unittest test_suite_name)
   set(_LLVM_TARGET_DEPENDENCIES
     "llvm_gtest_main;llvm_gtest;LLVMTestingSupport;LLVMSupport;LLVMDemangle")
 
+  if (NOT SYCL_COMPAT_UNITTEST_GCC_TOOLCHAIN STREQUAL "")
+    set(_GCC_TOOLCHAIN "--gcc-toolchain=${SYCL_COMPAT_UNITTEST_GCC_TOOLCHAIN}")
+  else()
+    set(_GCC_TOOLCHAIN "")
+  endif()
+
   foreach(_lib ${_LLVM_TARGET_DEPENDENCIES})
     list(APPEND _LIBRARIES $<TARGET_LINKER_FILE:${_lib}>)
   endforeach()
@@ -53,6 +60,7 @@ macro(add_sycl_library_unittest test_suite_name)
 
   add_custom_target(${_BIN_TARGET}
     COMMAND ${DEVICE_COMPILER_EXECUTABLE} -fsycl ${ARG_SOURCES}
+      ${_GCC_TOOLCHAIN}
       -o ${_OUTPUT_BIN}
       ${ARG_SYCL_EXTRA_FLAGS}
       ${_INTERNAL_EXTRA_FLAGS}

sycl/cmake/modules/AddSYCLLibraryUnitTest.cmake

@@ -566,7 +566,7 @@ class annotated_ref {
     annotated_ref(const annotated_ref&) = default;
     operator T() const;
     annotated_ref& operator=(const T &);
-    annotated_ref& operator=(const annotated_ref&) = default;
+    annotated_ref& operator=(const annotated_ref&);
   };
 } // namespace sycl::ext::oneapi::experimental
 ```
@@ -610,10 +610,11 @@ applying the annotations when the object is stored to memory.
 a|
 [source,c++]
 ----
-annotated_ref& operator=(const annotated_ref&) = default;
+annotated_ref& operator=(const annotated_ref&);
 ----
 |
-Assign from another `annotated_ref` object.
+Copy an object of type `T` from another `annotated_ref<T, ...>` object.
+Applying the annotations when the object is loaded from and stored to memory.
 
 |===
 

sycl/doc/extensions/experimental/sycl_ext_oneapi_annotated_ptr.asciidoc

@@ -371,15 +371,6 @@ size_t getLinearIndex(const T<Dims> &Index, const U<Dims> &Range) {
   return LinearIndex;
 }
 
-// Kernel set ID, used to group kernels (represented by OSModule & kernel name
-// pairs) into disjoint sets based on the kernel distribution among device
-// images.
-using KernelSetId = size_t;
-// Kernel set ID for kernels contained within the SPIR-V file specified via
-// environment.
-constexpr KernelSetId SpvFileKSId = 0;
-constexpr KernelSetId LastKSId = SpvFileKSId;
-
 template <typename T> struct InlineVariableHelper {
   static constexpr T value{};
 };

sycl/include/sycl/detail/common.hpp

@@ -107,6 +107,7 @@ class annotated_ref {
 template <typename T, typename... Props>
 class annotated_ref<T, detail::properties_t<Props...>> {
   using property_list_t = detail::properties_t<Props...>;
+  using this_t = annotated_ref<T, detail::properties_t<Props...>>;
 
 private:
   T *m_Ptr;
@@ -125,7 +126,7 @@ class annotated_ref<T, detail::properties_t<Props...>> {
 #endif
   }
 
-  annotated_ref &operator=(const T &Obj) {
+  this_t &operator=(const T &Obj) {
 #ifdef __SYCL_DEVICE_ONLY__
     *__builtin_intel_sycl_ptr_annotation(
         m_Ptr, detail::PropertyMetaInfo<Props>::name...,
@@ -136,7 +137,19 @@ class annotated_ref<T, detail::properties_t<Props...>> {
     return *this;
   }
 
-  annotated_ref &operator=(const annotated_ref &) = default;
+  this_t &operator=(const this_t &Obj) {
+    const T &t = Obj;
+    this->operator=(t);
+    return *this;
+  }
+
+  template <typename... OtherProperties>
+  this_t &operator=(
+      const annotated_ref<T, detail::properties_t<OtherProperties...>> &Obj) {
+    const T &t = Obj;
+    this->operator=(t);
+    return *this;
+  }
 
   PROPAGATE_OP(+=)
   PROPAGATE_OP(-=)

sycl/include/sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp

@@ -46,7 +46,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urContextGetInfo(
     size_t propSize, void *pContextInfo, size_t *pPropSizeRet) {
   UrReturnHelper ReturnValue(propSize, pContextInfo, pPropSizeRet);
 
-  switch (uint32_t{ContextInfoType}) {
+  switch (static_cast<uint32_t>(ContextInfoType)) {
   case UR_CONTEXT_INFO_NUM_DEVICES:
     return ReturnValue(1);
   case UR_CONTEXT_INFO_DEVICES:

sycl/plugins/unified_runtime/ur/adapters/cuda/context.cpp

@@ -6,6 +6,7 @@
 //
 //===-----------------------------------------------------------------===//
 
+#include <array>
 #include <cassert>
 #include <sstream>
 

sycl/plugins/unified_runtime/ur/adapters/cuda/device.cpp

@@ -71,10 +71,13 @@ urToCudaImageChannelFormat(ur_image_channel_type_t image_channel_type,
     PixelTypeSizeBytes = SIZE;                                                 \
     break;                                                                     \
   }
+// These new formats were brought in in CUDA 11.5
+#if CUDA_VERSION >= 11050
     CASE(UR_IMAGE_CHANNEL_TYPE_UNORM_INT8, CU_AD_FORMAT_UNORM_INT8X1, 1)
+    CASE(UR_IMAGE_CHANNEL_TYPE_UNORM_INT16, CU_AD_FORMAT_UNORM_INT16X1, 2)
+#endif
     CASE(UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8, CU_AD_FORMAT_UNSIGNED_INT8, 1)
     CASE(UR_IMAGE_CHANNEL_TYPE_SIGNED_INT8, CU_AD_FORMAT_SIGNED_INT8, 1)
-    CASE(UR_IMAGE_CHANNEL_TYPE_UNORM_INT16, CU_AD_FORMAT_UNORM_INT16X1, 2)
     CASE(UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16, CU_AD_FORMAT_UNSIGNED_INT16, 2)
     CASE(UR_IMAGE_CHANNEL_TYPE_SIGNED_INT16, CU_AD_FORMAT_SIGNED_INT16, 2)
     CASE(UR_IMAGE_CHANNEL_TYPE_HALF_FLOAT, CU_AD_FORMAT_HALF, 2)

sycl/plugins/unified_runtime/ur/adapters/cuda/image.cpp

@@ -10,6 +10,7 @@
 #include <cuda.h>
 #include <ur_api.h>
 
+#include <array>
 #include <atomic>
 #include <cassert>
 #include <numeric>

sycl/plugins/unified_runtime/ur/adapters/cuda/kernel.hpp

@@ -9,6 +9,7 @@
 
 #include <ur/ur.hpp>
 
+#include <algorithm>
 #include <cuda.h>
 #include <vector>
 

sycl/plugins/unified_runtime/ur/adapters/cuda/queue.hpp

@@ -14,6 +14,7 @@
 #include <mutex>
 #include <optional>
 #include <set>
+#include <unordered_set>
 
 #include <detail/pi_utils.hpp>
 #include <sycl/detail/defines_elementary.hpp>
@@ -51,11 +52,11 @@ struct DeviceGlobalMapEntry {
   std::string MUniqueId;
   // Pointer to the device_global on host.
   const void *MDeviceGlobalPtr = nullptr;
+  // Images device_global are used by.
+  std::unordered_set<RTDeviceBinaryImage *> MImages;
   // The image identifiers for the images using the device_global used by in the
   // cache.
   std::set<std::uintptr_t> MImageIdentifiers;
-  // The kernel-set IDs for the images using the device_global.
-  std::set<KernelSetId> MKSIds;
   // Size of the underlying type in the device_global.
   std::uint32_t MDeviceGlobalTSize = 0;
   // True if the device_global has been decorated with device_image_scope.
@@ -68,10 +69,11 @@ struct DeviceGlobalMapEntry {
 
   // Constructor for only initializing ID, type size, and device image scope
   // flag. The pointer to the device global will be initialized later.
-  DeviceGlobalMapEntry(std::string UniqueId, std::uintptr_t ImgId,
-                       KernelSetId KSId, std::uint32_t DeviceGlobalTSize,
+  DeviceGlobalMapEntry(std::string UniqueId, RTDeviceBinaryImage *Img,
+                       std::uint32_t DeviceGlobalTSize,
                        bool IsDeviceImageScopeDecorated)
-      : MUniqueId(UniqueId), MImageIdentifiers{ImgId}, MKSIds{KSId},
+      : MUniqueId(UniqueId), MImages{Img},
+        MImageIdentifiers{reinterpret_cast<uintptr_t>(Img)},
         MDeviceGlobalTSize(DeviceGlobalTSize),
         MIsDeviceImageScopeDecorated(IsDeviceImageScopeDecorated) {}
 
@@ -85,8 +87,7 @@ struct DeviceGlobalMapEntry {
 
   // Initialize the device_global's element type size and the flag signalling
   // if the device_global has the device_image_scope property.
-  void initialize(std::uintptr_t ImgId, KernelSetId KSId,
-                  std::uint32_t DeviceGlobalTSize,
+  void initialize(RTDeviceBinaryImage *Img, std::uint32_t DeviceGlobalTSize,
                   bool IsDeviceImageScopeDecorated) {
     if (MDeviceGlobalTSize != 0) {
       // The device global entry has already been initialized. This can happen
@@ -99,8 +100,8 @@ struct DeviceGlobalMapEntry {
           "Device global intializations disagree on image scope decoration.");
       return;
     }
-    MImageIdentifiers.insert(ImgId);
-    MKSIds.insert(KSId);
+    MImages.insert(Img);
+    MImageIdentifiers.insert(reinterpret_cast<uintptr_t>(Img));
     MDeviceGlobalTSize = DeviceGlobalTSize;
     MIsDeviceImageScopeDecorated = IsDeviceImageScopeDecorated;
   }

sycl/source/detail/device_global_map_entry.hpp

@@ -1133,13 +1133,13 @@ getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue,
   assert(DeviceGlobalEntry->MIsDeviceImageScopeDecorated &&
          "device_global is not device image scope decorated.");
 
-  // If the device global is used in multiple kernel sets we cannot proceed.
-  if (DeviceGlobalEntry->MKSIds.size() > 1)
+  // If the device global is used in multiple device images we cannot proceed.
+  if (DeviceGlobalEntry->MImageIdentifiers.size() > 1)
     throw sycl::exception(make_error_code(errc::invalid),
                           "More than one image exists with the device_global.");
 
   // If there are no kernels using the device_global we cannot proceed.
-  if (DeviceGlobalEntry->MKSIds.size() == 0)
+  if (DeviceGlobalEntry->MImageIdentifiers.size() == 0)
     throw sycl::exception(make_error_code(errc::invalid),
                           "No image exists with the device_global.");
 
@@ -1153,9 +1153,9 @@ getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue,
 
   // If there was no cached program, build one.
   auto Context = createSyclObjFromImpl<context>(ContextImpl);
-  KernelSetId KSId = *DeviceGlobalEntry->MKSIds.begin();
   ProgramManager &PM = ProgramManager::getInstance();
-  RTDeviceBinaryImage &Img = PM.getDeviceImage(KSId, Context, Device);
+  RTDeviceBinaryImage &Img =
+      PM.getDeviceImage(DeviceGlobalEntry->MImages, Context, Device);
   device_image_plain DeviceImage =
       PM.getDeviceImageFromBinaryImage(&Img, Context, Device);
   device_image_plain BuiltImage = PM.build(DeviceImage, {Device}, {});