Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/workflows/windows_test_comment_trigger.yml

@@ -28,7 +28,7 @@ jobs:
           const pr = await github.rest.pulls.get({
             owner: context.issue.owner,
             repo: context.issue.repo,
-            pull_number: number
+            pull_number: context.issue.number
           });
           return pr.data.head.sha
     - name: update_pr_status_pending

libclc/generic/libspirv/integer/ctz.cl

@@ -34,15 +34,15 @@ _CLC_OVERLOAD _CLC_DEF ulong __spirv_ocl_ctz(ulong x) {
 }
 
 _CLC_OVERLOAD _CLC_DEF char __spirv_ocl_ctz(char x) {
-  return __spirv_ocl_ctz((ushort)(uchar)x) - 8;
+  return x ? __spirv_ocl_ctz((ushort)(uchar)x) : 8;
 }
 
 _CLC_OVERLOAD _CLC_DEF schar __spirv_ocl_ctz(schar x) {
-  return __spirv_ocl_ctz((ushort)(uchar)x) - 8;
+  return x ? __spirv_ocl_ctz((ushort)(uchar)x) : 8;
 }
 
 _CLC_OVERLOAD _CLC_DEF uchar __spirv_ocl_ctz(uchar x) {
-  return __spirv_ocl_ctz((ushort)x) - 8;
+  return x ? __spirv_ocl_ctz((ushort)x) : 8;
 }
 
 _CLC_UNARY_VECTORIZE(_CLC_OVERLOAD _CLC_DEF, char, __spirv_ocl_ctz, char)

llvm/lib/Passes/PassBuilderPipelines.cpp

@@ -635,9 +635,7 @@ PassBuilder::buildFunctionSimplificationPipeline(OptimizationLevel Level,
 
   // Try vectorization/scalarization transforms that are both improvements
   // themselves and can allow further folds with GVN and InstCombine.
-  // Disable for SYCL until SPIR-V reader is updated for all drivers.
-  if (!SYCLOptimizationMode)
-    FPM.addPass(VectorCombinePass(/*TryEarlyFoldsOnly=*/true));
+  FPM.addPass(VectorCombinePass(/*TryEarlyFoldsOnly=*/true));
 
   // Eliminate redundancies.
   FPM.addPass(MergedLoadStoreMotionPass());

llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp

@@ -1195,6 +1195,25 @@ translateSpirvGlobalUses(LoadInst *LI, StringRef SpirvGlobalName,
     return;
   }
 
+  // As an optimization of accesses of the first element for vector SPIRV
+  // globals sometimes the load will return a scalar and the uses can be of any
+  // pattern. In this case, generate GenX calls to access the first element and
+  // update the use to instead use the GenX call result rather than the load
+  // result.
+  if (!LI->getType()->isVectorTy()) {
+    // Copy users to seperate container for safe modification
+    // during iteration.
+    SmallVector<User *> Users(LI->users());
+    for (User *LU : Users) {
+      Instruction *Inst = cast<Instruction>(LU);
+      NewInst =
+          generateSpirvGlobalGenX(Inst, SpirvGlobalName, /*IndexValue=*/0);
+      LU->replaceUsesOfWith(LI, NewInst);
+    }
+    InstsToErase.push_back(LI);
+    return;
+  }
+
   // Only loads from _vector_ SPIRV globals reach here now. Their users are
   // expected to be ExtractElementInst only, and they are
   // replaced in this loop. When loads from _scalar_ SPIRV globals are handled

llvm/test/SYCLLowerIR/ESIMD/lower_spirv_intrins_vector_treated_as_scalar.ll

@@ -0,0 +1,32 @@
+; RUN: opt -opaque-pointers < %s -passes=LowerESIMD -S | FileCheck %s
+
+; This test checks we lower vector SPIRV globals correctly if
+; it is accessed as a scalar as an optimization to get the first element
+
+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 dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+
+define spir_kernel void @"__spirv_GlobalInvocationId_xyz"(i64 addrspace(1)* %_arg_) {
+; CHECK-LABEL: @__spirv_GlobalInvocationId_xyz(
+; CHECK-NEXT:  entry:
+; CHECK-NEXT:    [[DOTESIMD6:%.*]] = call <3 x i32> @llvm.genx.local.id.v3i32()
+; CHECK-NEXT:    [[LOCAL_ID_X:%.*]] = extractelement <3 x i32> [[DOTESIMD6]], i32 0
+; CHECK-NEXT:    [[LOCAL_ID_X_CAST_TY:%.*]] = zext i32 [[LOCAL_ID_X]] to i64
+; CHECK-NEXT:    [[DOTESIMD7:%.*]] = call <3 x i32> @llvm.genx.local.size.v3i32()
+; CHECK-NEXT:    [[WGSIZE_X:%.*]] = extractelement <3 x i32> [[DOTESIMD7]], i32 0
+; CHECK-NEXT:    [[WGSIZE_X_CAST_TY:%.*]] = zext i32 [[WGSIZE_X]] to i64
+; CHECK-NEXT:    [[GROUP_ID_X:%.*]] = call i32 @llvm.genx.group.id.x()
+; CHECK-NEXT:    [[GROUP_ID_X_CAST_TY:%.*]] = zext i32 [[GROUP_ID_X]] to i64
+; CHECK-NEXT:    [[MUL8:%.*]] = mul i64 [[WGSIZE_X_CAST_TY]], [[GROUP_ID_X_CAST_TY]]
+; CHECK-NEXT:    [[ADD9:%.*]] = add i64 [[LOCAL_ID_X_CAST_TY]], [[MUL8]]
+; CHECK-NEXT:    [[MUL10:%.*]] = shl nuw nsw i64 [[ADD9]], 5
+
+; Verify that the attribute is deleted from GenX declaration
+; CHECK-NOT: readnone
+entry:
+ %0 = load i64, ptr addrspace(1) @__spirv_BuiltInGlobalInvocationId, align 32
+ %mul.i = shl nuw nsw i64 %0, 5
+ ret void
+}

sycl/cmake/modules/AddBoostMp11Headers.cmake

@@ -45,7 +45,7 @@ add_custom_target(boost_mp11-headers
 # ${BOOST_MP11_DESTINATION_DIR}
 add_custom_command(
   OUTPUT ${OUT_HEADERS_BOOST_MP11}
-  DEPENDS ${HEADERS_BOOST_MP11}
+  DEPENDS ${HEADERS_BOOST_MP11} ${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/PreprocessBoostMp11Headers.cmake
   COMMAND ${CMAKE_COMMAND}
     -DIN=${BOOST_MP11_SOURCE_DIR}/include/boost
     -DOUT=${BOOST_MP11_DESTINATION_DIR}

sycl/cmake/modules/PreprocessBoostMp11Headers.cmake

@@ -37,12 +37,12 @@ function(preprocess_mp11_header)
   # `namespace sycl { namespace detail { namespace boost { ... } } }`
   string(REGEX REPLACE
     "(\n[ \t]*namespace[ \t\n\r]+boost)"
-    "namespace sycl\n{\nnamespace detail\n{\\1"
+    "namespace sycl\n{\ninline namespace _V1\n{\nnamespace detail\n{\\1"
     FILE_CONTENTS "${FILE_CONTENTS}")
   # ... use '} // namespace boost' as a marker for end-of-scope '}' replacement
   string(REGEX REPLACE
     "(\n[ \t]*}[ \t]*//[ \t]*namespace[ \t]+boost[ \t]*\n)"
-    "\\1} // namespace detail\n} // namespace sycl\n"
+    "\\1} // namespace detail\n} // namespace _V1\n} // namespace sycl\n"
     FILE_CONTENTS "${FILE_CONTENTS}")
   # 3) replace `boost` in `#include <boost/...>` or `#include "boost/..."` with
   # `sycl/detail/boost`

sycl/include/sycl/ext/oneapi/properties/property_utils.hpp

@@ -10,6 +10,7 @@
 
 #include <sycl/detail/property_helper.hpp>
 #include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/detail/boost/mp11.hpp>
 
 #include <tuple>
 
@@ -106,93 +107,20 @@ template <typename RHS> struct SelectNonVoid<void, RHS> {
   using type = RHS;
 };
 
-// Merges two tuples by recursively extracting the type with the minimum
-// PropertyID in the two tuples and prepending it to the merging of the
-// remaining elements.
-template <typename T1, typename T2> struct Merge {};
-template <typename... LTs> struct Merge<std::tuple<LTs...>, std::tuple<>> {
-  using type = std::tuple<LTs...>;
+// Sort types accoring to their PropertyID.
+struct SortByPropertyId {
+  template <typename T1, typename T2>
+  using fn = sycl::detail::boost::mp11::mp_bool<(PropertyID<T1>::value <
+                                                 PropertyID<T2>::value)>;
 };
-template <typename... RTs> struct Merge<std::tuple<>, std::tuple<RTs...>> {
-  using type = std::tuple<RTs...>;
-};
-template <typename... LTs, typename... RTs>
-struct Merge<std::tuple<LTs...>, std::tuple<RTs...>> {
-  using l_head = GetFirstType<LTs...>;
-  using r_head = GetFirstType<RTs...>;
-  static constexpr bool left_has_min =
-      PropertyID<l_head>::value < PropertyID<r_head>::value;
-  using l_split = HeadSplit<std::tuple<LTs...>, left_has_min>;
-  using r_split = HeadSplit<std::tuple<RTs...>, !left_has_min>;
-  using min = typename SelectNonVoid<typename l_split::htype,
-                                     typename r_split::htype>::type;
-  using merge_tails =
-      typename Merge<typename l_split::ttype, typename r_split::ttype>::type;
-  using type = typename PrependTuple<min, merge_tails>::type;
-};
-
-// Creates pairs of tuples with a single element from a tuple with N elements.
-// Resulting tuple will have ceil(N/2) elements.
-template <typename...> struct CreateTuplePairs {
-  using type = typename std::tuple<>;
-};
-template <typename T> struct CreateTuplePairs<T> {
-  using type = typename std::tuple<std::pair<std::tuple<T>, std::tuple<>>>;
-};
-template <typename L, typename R, typename... Rest>
-struct CreateTuplePairs<L, R, Rest...> {
-  using type =
-      typename PrependTuple<std::pair<std::tuple<L>, std::tuple<R>>,
-                            typename CreateTuplePairs<Rest...>::type>::type;
-};
-
-// Merges pairs of tuples and creates new pairs of the merged pairs. Let N be
-// the number of pairs in the supplied tuple, then the resulting tuple will
-// contain ceil(N/2) pairs of tuples.
-template <typename T> struct MergePairs {
-  using type = std::tuple<>;
-};
-template <typename... LTs, typename... RTs, typename... Rest>
-struct MergePairs<
-    std::tuple<std::pair<std::tuple<LTs...>, std::tuple<RTs...>>, Rest...>> {
-  using merged = typename Merge<std::tuple<LTs...>, std::tuple<RTs...>>::type;
-  using type = std::tuple<std::pair<merged, std::tuple<>>>;
-};
-template <typename... LLTs, typename... LRTs, typename... RLTs,
-          typename... RRTs, typename... Rest>
-struct MergePairs<
-    std::tuple<std::pair<std::tuple<LLTs...>, std::tuple<LRTs...>>,
-               std::pair<std::tuple<RLTs...>, std::tuple<RRTs...>>, Rest...>> {
-  using lmerged =
-      typename Merge<std::tuple<LLTs...>, std::tuple<LRTs...>>::type;
-  using rmerged =
-      typename Merge<std::tuple<RLTs...>, std::tuple<RRTs...>>::type;
-  using type = typename PrependTuple<
-      std::pair<lmerged, rmerged>,
-      typename MergePairs<std::tuple<Rest...>>::type>::type;
-};
-
-// Recursively merges all pairs of tuples until only a single pair of tuples
-// is left, where the right element of the pair is an empty tuple.
-template <typename T> struct MergeAll {};
-template <typename... Ts> struct MergeAll<std::tuple<Ts...>> {
-  using type = std::tuple<Ts...>;
-};
-template <typename... Ts>
-struct MergeAll<std::tuple<std::pair<std::tuple<Ts...>, std::tuple<>>>> {
-  using type = std::tuple<Ts...>;
-};
-template <typename T, typename... Ts> struct MergeAll<std::tuple<T, Ts...>> {
-  using reduced = typename MergePairs<std::tuple<T, Ts...>>::type;
-  using type = typename MergeAll<reduced>::type;
-};
-
-// Performs merge-sort on types with PropertyID.
 template <typename... Ts> struct Sorted {
   static_assert(detail::AllPropertyValues<std::tuple<Ts...>>::value,
                 "Unrecognized property in property list.");
-  using split = typename CreateTuplePairs<Ts...>::type;
-  using type = typename MergeAll<split>::type;
+  using properties = sycl::detail::boost::mp11::mp_list<Ts...>;
+  using sortedProperties =
+      sycl::detail::boost::mp11::mp_sort_q<properties, SortByPropertyId>;
+  using type =
+      sycl::detail::boost::mp11::mp_rename<sortedProperties, std::tuple>;
 };
 
 // Checks if the types in a tuple are sorted w.r.t. their PropertyID.

sycl/include/sycl/marray.hpp

@@ -72,17 +72,25 @@ template <typename DataT, typename ArgT, typename... ArgTN>
 struct ArrayCreator<DataT, ArgT, ArgTN...> {
   static constexpr std::array<DataT, GetMArrayArgsSize<ArgT, ArgTN...>::value>
   Create(const ArgT &Arg, const ArgTN &...Args) {
-    return ConcatArrays(std::array<DataT, 1>{static_cast<DataT>(Arg)},
-                        ArrayCreator<DataT, ArgTN...>::Create(Args...));
+    std::array<DataT, 1> ImmArray{static_cast<DataT>(Arg)};
+    if constexpr (sizeof...(Args))
+      return ConcatArrays(ImmArray,
+                          ArrayCreator<DataT, ArgTN...>::Create(Args...));
+    else
+      return ImmArray;
   }
 };
 template <typename DataT, typename T, std::size_t N, typename... ArgTN>
 struct ArrayCreator<DataT, marray<T, N>, ArgTN...> {
   static constexpr std::array<DataT,
                               GetMArrayArgsSize<marray<T, N>, ArgTN...>::value>
   Create(const marray<T, N> &Arg, const ArgTN &...Args) {
-    return ConcatArrays(MArrayToArray<DataT>(Arg),
-                        ArrayCreator<DataT, ArgTN...>::Create(Args...));
+    auto ImmArray = MArrayToArray<DataT>(Arg);
+    if constexpr (sizeof...(Args))
+      return ConcatArrays(ImmArray,
+                          ArrayCreator<DataT, ArgTN...>::Create(Args...));
+    else
+      return ImmArray;
   }
 };
 template <typename DataT> struct ArrayCreator<DataT> {

sycl/plugins/cuda/pi_cuda.cpp

@@ -1966,6 +1966,53 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
                         name.data());
   }
 
+  case PI_DEVICE_INFO_MAX_MEM_BANDWIDTH: {
+    int major = 0;
+    sycl::detail::pi::assertion(
+        cuDeviceGetAttribute(&major,
+                             CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,
+                             device->get()) == CUDA_SUCCESS);
+
+    int minor = 0;
+    sycl::detail::pi::assertion(
+        cuDeviceGetAttribute(&minor,
+                             CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,
+                             device->get()) == CUDA_SUCCESS);
+
+    // Some specific devices seem to need special handling. See reference
+    // https://github.com/jeffhammond/HPCInfo/blob/master/cuda/gpu-detect.cu
+    bool is_xavier_agx = major == 7 && minor == 2;
+    bool is_orin_agx = major == 8 && minor == 7;
+
+    int memory_clock_khz = 0;
+    if (is_xavier_agx) {
+      memory_clock_khz = 2133000;
+    } else if (is_orin_agx) {
+      memory_clock_khz = 3200000;
+    } else {
+      sycl::detail::pi::assertion(
+          cuDeviceGetAttribute(&memory_clock_khz,
+                               CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE,
+                               device->get()) == CUDA_SUCCESS);
+    }
+
+    int memory_bus_width = 0;
+    if (is_orin_agx) {
+      memory_bus_width = 256;
+    } else {
+      sycl::detail::pi::assertion(
+          cuDeviceGetAttribute(&memory_bus_width,
+                               CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH,
+                               device->get()) == CUDA_SUCCESS);
+    }
+
+    uint64_t memory_bandwidth =
+        uint64_t(memory_clock_khz) * memory_bus_width * 250;
+
+    return getInfo(param_value_size, param_value, param_value_size_ret,
+                   memory_bandwidth);
+  }
+
     // TODO: Investigate if this information is available on CUDA.
   case PI_DEVICE_INFO_PCI_ADDRESS:
   case PI_DEVICE_INFO_GPU_EU_COUNT:
@@ -1974,7 +2021,6 @@ pi_result cuda_piDeviceGetInfo(pi_device device, pi_device_info param_name,
   case PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE:
   case PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE:
   case PI_DEVICE_INFO_GPU_HW_THREADS_PER_EU:
-  case PI_DEVICE_INFO_MAX_MEM_BANDWIDTH:
     return PI_ERROR_INVALID_VALUE;
 
   default:
@@ -2985,7 +3031,7 @@ pi_result cuda_piEnqueueKernelLaunch(
             return PI_ERROR_INVALID_WORK_GROUP_SIZE;
 
           if (local_work_size[dim] > maxThreadsPerBlock[dim])
-            return PI_ERROR_INVALID_WORK_ITEM_SIZE;
+            return PI_ERROR_INVALID_WORK_GROUP_SIZE;
           // Checks that local work sizes are a divisor of the global work sizes
           // which includes that the local work sizes are neither larger than
           // the global work sizes and not 0.

sycl/plugins/hip/pi_hip.cpp

@@ -2867,7 +2867,7 @@ pi_result hip_piEnqueueKernelLaunch(
     if (providedLocalWorkGroupSize) {
       auto isValid = [&](int dim) {
         if (local_work_size[dim] > maxThreadsPerBlock[dim])
-          return PI_ERROR_INVALID_WORK_ITEM_SIZE;
+          return PI_ERROR_INVALID_WORK_GROUP_SIZE;
         // Checks that local work sizes are a divisor of the global work sizes
         // which includes that the local work sizes are neither larger than the
         // global work sizes and not 0.

sycl/source/detail/context_impl.cpp

@@ -339,7 +339,7 @@ std::vector<RT::PiEvent> context_impl::initializeDeviceGlobals(
       // Write the pointer to the device global and store the event in the
       // initialize events list.
       RT::PiEvent InitEvent;
-      void *USMPtr = DeviceGlobalUSM.getPtr();
+      void *const &USMPtr = DeviceGlobalUSM.getPtr();
       Plugin.call<PiApiKind::piextEnqueueDeviceGlobalVariableWrite>(
           QueueImpl->getHandleRef(), NativePrg,
           DeviceGlobalEntry->MUniqueId.c_str(), false, sizeof(void *), 0,

sycl/source/detail/device_global_map_entry.hpp

@@ -64,7 +64,7 @@ struct DeviceGlobalUSMMem {
   DeviceGlobalUSMMem(void *Ptr) : MPtr(Ptr) {}
   ~DeviceGlobalUSMMem();
 
-  void *getPtr() const noexcept { return MPtr; }
+  void *const &getPtr() const noexcept { return MPtr; }
 
   // Gets the zero-initialization event if it exists. If not the OwnedPiEvent
   // will contain no event.

sycl/test-e2e/AOT/cpu.cpp

@@ -0,0 +1,15 @@
+//==--- cpu.cpp - AOT compilation for cpu devices using opencl-aot --------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: opencl-aot, cpu
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 %S/Inputs/aot.cpp -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+
+// Test that opencl-aot can handle multiple build options.
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64  %S/Inputs/aot.cpp -Xsycl-target-backend "--bo=-g" -Xsycl-target-backend "--bo=-cl-opt-disable" -o %t2.out

sycl/test-e2e/AOT/gpu.cpp

@@ -0,0 +1,14 @@
+//==--- gpu.cpp - AOT compilation for gen devices using GEN compiler ------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: ocloc, gpu
+// UNSUPPORTED: cuda
+// CUDA is not compatible with SPIR.
+//
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %S/Inputs/aot.cpp -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out