Merge from 'sycl' to 'sycl-web' (4 commits)

Author: iclsrc

devops/dependencies.json

@@ -19,9 +19,9 @@
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "level_zero": {
-      "github_tag": "v1.15.8",
-      "version": "v1.15.8",
-      "url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.15.8",
+      "github_tag": "v1.15.13",
+      "version": "v1.15.13",
+      "url": "https://github.com/oneapi-src/level-zero/releases/tag/v1.15.13",
       "root": "{DEPS_ROOT}/opencl/runtime/linux/oclgpu"
     },
     "tbb": {

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -405,9 +405,7 @@ static void copyBetweenPrivateAndShadow(Value *L, GlobalVariable *Shadow,
   assert(T && "Unexpected type");
 
   if (T->isAggregateType()) {
-    // TODO: we should use methods which directly return MaybeAlign once such
-    // are added to LLVM for AllocaInst and GlobalVariable
-    auto ShdAlign = MaybeAlign(Shadow->getAlignment());
+    auto ShdAlign = Shadow->getAlign();
     Module &M = *Shadow->getParent();
     auto SizeVal = M.getDataLayout().getTypeStoreSize(T);
     auto Size = ConstantInt::get(getSizeTTy(M), SizeVal);

sycl/include/sycl/builtins_preview.hpp

@@ -170,16 +170,22 @@ auto builtin_default_host_impl(FuncTy F, const Ts &...x) {
 template <typename FuncTy, typename... Ts>
 auto builtin_delegate_to_scalar(FuncTy F, const Ts &...x) {
   using T = typename first_type<Ts...>::type;
-  if constexpr (is_vec_or_swizzle_v<T>) {
-    using ret_elem_type = decltype(F(x[0]...));
-    // TODO: using r{} to avoid Werror. Not sure if ok.
-    vec<ret_elem_type, T::size()> r{};
-    loop<T::size()>([&](auto idx) { r[idx] = F(x[idx]...); });
-    return r;
+  static_assert(is_vec_or_swizzle_v<T> || is_marray_v<T>);
+
+  constexpr auto Size = T::size();
+  using ret_elem_type = decltype(F(x[0]...));
+  std::conditional_t<is_marray_v<T>, marray<ret_elem_type, Size>,
+                     vec<ret_elem_type, Size>>
+      r{};
+
+  if constexpr (is_marray_v<T>) {
+    for (size_t i = 0; i < Size; ++i)
+      r[i] = F(x[i]...);
   } else {
-    static_assert(is_marray_v<T>);
-    return builtin_marray_impl(F, x...);
+    loop<Size>([&](auto idx) { r[idx] = F(x[idx]...); });
   }
+
+  return r;
 }
 
 template <typename T>

sycl/test-e2e/Basic/built-ins/marray_integer.cpp

@@ -0,0 +1,49 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+// RUN: %if preview-breaking-changes-supported %{ %{build} -fpreview-breaking-changes -o %t_preview.out %}
+// RUN: %if preview-breaking-changes-supported %{ %{run} %t_preview.out%}
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  using namespace sycl;
+  queue q;
+
+  auto Test = [&](auto F, auto Expected, auto... Args) {
+    std::tuple ArgsTuple{Args...};
+    auto Result = std::apply(F, ArgsTuple);
+    static_assert(std::is_same_v<decltype(Expected), decltype(Result)>);
+
+    auto Equal = [](auto x, auto y) {
+      for (size_t i = 0; i < x.size(); ++i)
+        if (x[i] != y[i])
+          return false;
+
+      return true;
+    };
+
+    assert(Equal(Result, Expected));
+
+    buffer<bool, 1> ResultBuf{1};
+    q.submit([&](handler &cgh) {
+      accessor Result{ResultBuf, cgh};
+      cgh.single_task([=]() {
+        auto R = std::apply(F, ArgsTuple);
+        static_assert(std::is_same_v<decltype(Expected), decltype(R)>);
+        Result[0] = Equal(R, Expected);
+      });
+    });
+    assert(host_accessor{ResultBuf}[0]);
+  };
+
+  {
+    // Test upsample:
+    auto Upsample = [](auto... xs) { return upsample(xs...); };
+    Test(Upsample, marray<int16_t, 2>{0x203, 0x302}, marray<int8_t, 2>{2, 3},
+         marray<uint8_t, 2>{3, 2});
+    Test(Upsample, marray<uint16_t, 2>{0x203, 0x302}, marray<uint8_t, 2>{2, 3},
+         marray<uint8_t, 2>{3, 2});
+  }
+
+  return 0;
+}

sycl/test-e2e/Matrix/SG32/joint_matrix_bfloat16_rowmajorA_rowmajorB.cpp renamed to sycl/test-e2e/Matrix/SG32/joint_matrix_rowmajorA_rowmajorB.cpp

@@ -1,27 +1,27 @@
-//==--joint_matrix_bfloat16_rowmajorA_rowmajorB.cpp  - DPC++ joint_matrix---==//
+//==--------joint_matrix_rowmajorA_rowmajorB.cpp  - DPC++ joint_matrix------==//
 //
 // 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: matrix
+// This tests support of row major layout for matrix B which does automatic VNNI
+// REQUIRES: aspect-ext_intel_matrix
 // REQUIRES-INTEL-DRIVER: lin: 27501, win: 101.4943
+// VNNI transform and sub-group size 32 are not supported yet on DG2 by IGC
+// UNSUPPORTED: gpu-intel-dg2
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// This tests support of row major layout for matrix B which does automatic VNNI
-// transform. This is currently only available on AMX
-
+// Sub-group size 32 support for this test is not currently available in IGC
 // XFAIL: gpu
 
 #include "../common.hpp"
 
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-constexpr size_t SG_SZ = 32;
-constexpr size_t TN = 16;
+#define SG_SZ 32
 
-#include "../joint_matrix_bfloat16_rowmajorA_rowmajorB_impl.hpp"
+#include "../joint_matrix_rowmajorA_rowmajorB_impl.hpp"

sycl/test-e2e/Matrix/joint_matrix_bfloat16_rowmajorA_rowmajorB_impl.hpp

@@ -1,11 +1,13 @@
-//==--joint_matrix_bfloat16_rowmajorA_rowmajorB.cpp  - DPC++ joint_matrix---==//
+//==-------joint_matrix_rowmajorA_rowmajorB.cpp  - DPC++ joint_matrix-------==//
 //
 // 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: matrix
+// REQUIRES: aspect-ext_intel_matrix
+// VNNI transform is not supported yet by IGC on DG2
+// UNSUPPORTED: gpu-intel-dg2
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
@@ -18,7 +20,4 @@
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
 
-#define SG_SZ 16
-constexpr size_t TN = 16;
-
-#include "joint_matrix_bfloat16_rowmajorA_rowmajorB_impl.hpp"
+#include "joint_matrix_rowmajorA_rowmajorB_impl.hpp"