From 1050ea809aa6f6a5c98bf6e52c7a2b5eb7786cdf Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Mon, 12 Feb 2024 14:04:20 -0800
Subject: [PATCH 1/2] [SYCL] Fix a bug with sycl::sample on marray types

---
 sycl/include/sycl/builtins_preview.hpp        | 22 ++++++---
 .../Basic/built-ins/marray_integer.cpp        | 49 +++++++++++++++++++
 2 files changed, 63 insertions(+), 8 deletions(-)
 create mode 100644 sycl/test-e2e/Basic/built-ins/marray_integer.cpp

diff --git a/sycl/include/sycl/builtins_preview.hpp b/sycl/include/sycl/builtins_preview.hpp
index 91fd50d2ec4f2..67f426c308e8a 100644
--- a/sycl/include/sycl/builtins_preview.hpp
+++ b/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 N = T::size();
+  using ret_elem_type = decltype(F(x[0]...));
+  std::conditional_t<is_marray_v<T>, marray<ret_elem_type, N>,
+                     vec<ret_elem_type, N>>
+      r{};
+
+  if constexpr (is_marray_v<T>) {
+    for (size_t i = 0; i < T::size(); ++i)
+      r[i] = F(x[i]...);
   } else {
-    static_assert(is_marray_v<T>);
-    return builtin_marray_impl(F, x...);
+    loop<T::size()>([&](auto idx) { r[idx] = F(x[idx]...); });
   }
+
+  return r;
 }
 
 template <typename T>
diff --git a/sycl/test-e2e/Basic/built-ins/marray_integer.cpp b/sycl/test-e2e/Basic/built-ins/marray_integer.cpp
new file mode 100644
index 0000000000000..3883c99f4fde9
--- /dev/null
+++ b/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;
+}

From b158e193c264dc803c38c97e2ec3b279e7ac4deb Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 13 Feb 2024 08:28:15 -0800
Subject: [PATCH 2/2] address review comments

---
 sycl/include/sycl/builtins_preview.hpp | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/sycl/include/sycl/builtins_preview.hpp b/sycl/include/sycl/builtins_preview.hpp
index 67f426c308e8a..03c1c401fbcd4 100644
--- a/sycl/include/sycl/builtins_preview.hpp
+++ b/sycl/include/sycl/builtins_preview.hpp
@@ -172,17 +172,17 @@ auto builtin_delegate_to_scalar(FuncTy F, const Ts &...x) {
   using T = typename first_type<Ts...>::type;
   static_assert(is_vec_or_swizzle_v<T> || is_marray_v<T>);
 
-  constexpr auto N = T::size();
+  constexpr auto Size = T::size();
   using ret_elem_type = decltype(F(x[0]...));
-  std::conditional_t<is_marray_v<T>, marray<ret_elem_type, N>,
-                     vec<ret_elem_type, N>>
+  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 < T::size(); ++i)
+    for (size_t i = 0; i < Size; ++i)
       r[i] = F(x[i]...);
   } else {
-    loop<T::size()>([&](auto idx) { r[idx] = F(x[idx]...); });
+    loop<Size>([&](auto idx) { r[idx] = F(x[idx]...); });
   }
 
   return r;