Skip to content

Commit 7e11e48

Browse files
authored
[SYCL][ESIMD] Fix logic and arithmetic operators for simd_view with scalar (#4680)
These operators should call corresponding operator for simd_obj_imp with scalar instead of constructing simd object from input scalar themselves. simd_obj_impl operators will do the necessary type promotion for operands if needed which is not done in the current implementation.
1 parent 6fed6b2 commit 7e11e48

File tree

2 files changed

+24
-4
lines changed

2 files changed

+24
-4
lines changed

sycl/include/sycl/ext/intel/experimental/esimd/detail/operators.hpp

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -290,8 +290,7 @@ namespace __SEIEE {
290290
class = std::enable_if_t< \
291291
__SEIEED::is_any_simd_view_type_v<SimdViewT2> && COND>> \
292292
inline auto operator BINOP(T1 LHS, const SimdViewT2 &RHS) { \
293-
using SimdT = typename SimdViewT2::value_type; \
294-
return SimdT(LHS) BINOP RHS.read(); \
293+
return LHS BINOP RHS.read(); \
295294
} \
296295
\
297296
/* simd_view BINOP SCALAR */ \
@@ -301,8 +300,7 @@ namespace __SEIEE {
301300
class = std::enable_if_t< \
302301
__SEIEED::is_any_simd_view_type_v<SimdViewT1> && COND>> \
303302
inline auto operator BINOP(const SimdViewT1 &LHS, T2 RHS) { \
304-
using SimdT = typename SimdViewT1::value_type; \
305-
return LHS.read() BINOP SimdT(RHS); \
303+
return LHS.read() BINOP RHS; \
306304
}
307305

308306
#define __ESIMD_BITWISE_OP_FILTER \

sycl/test/esimd/simd_view_bin_op.cpp

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
// RUN: %clangxx -fsycl -fsycl-device-only -S -emit-llvm -x c++ %s -o - | FileCheck %s
2+
3+
// This test checks that arithmetic opration on simd_view and scalar does not
4+
// truncate scalar to the view's element type.
5+
6+
#include <sycl/ext/intel/experimental/esimd.hpp>
7+
8+
using namespace sycl::ext::intel::experimental::esimd;
9+
10+
simd<short, 4> test1(int X, simd<short, 16> Y) __attribute__((sycl_device)) {
11+
// CHECK-LABEL: test1
12+
// CHECK-SAME: i32 [[X:%.+]],
13+
// CHECK-NOT: trunc i32 [[X]] to i16
14+
return X * Y.select<4, 1>(0);
15+
}
16+
17+
simd<short, 4> test2(int X, simd<short, 16> Y) __attribute__((sycl_device)) {
18+
// CHECK-LABEL: test2
19+
// CHECK-SAME: i32 [[X:%.+]],
20+
// CHECK-NOT: trunc i32 [[X]] to i16
21+
return Y.select<4, 1>(0) * X;
22+
}

0 commit comments

Comments
 (0)