[SYCL] Fix a bug with sycl::sample on marray types (#12694)

Author: aelovikov-intel

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;
+}