From 545745352a696c5399a62242c78e0dbf618aac3e Mon Sep 17 00:00:00 2001 From: Konstantin S Bobrovsky Date: Tue, 28 Dec 2021 16:02:35 -0800 Subject: [PATCH 1/2] [ESIMD] Add basic operations code generation test for half type. Signed-off-by: Konstantin S Bobrovsky --- sycl/test/esimd/sycl_half_basic_ops.cpp | 87 +++++++++++++++++++++++++ 1 file changed, 87 insertions(+) create mode 100644 sycl/test/esimd/sycl_half_basic_ops.cpp diff --git a/sycl/test/esimd/sycl_half_basic_ops.cpp b/sycl/test/esimd/sycl_half_basic_ops.cpp new file mode 100644 index 0000000000000..c74c5e4f23be9 --- /dev/null +++ b/sycl/test/esimd/sycl_half_basic_ops.cpp @@ -0,0 +1,87 @@ +// RUN: %clangxx -fsycl -fsycl-device-only -S %s -o %t +// RUN: sycl-post-link -split-esimd -lower-esimd -S %t -o %t.table +// RUN: FileCheck %s -input-file=%t_esimd_0.ll + +// The test checks that there are no unexpected extra conversions or intrinsic +// calls added by the API headers or compiler when generating code +// for basic C++ operations on simd values. + +#include + +using namespace sycl::ext::intel::experimental::esimd; +using namespace sycl::ext::intel::experimental; +using namespace sycl; + +// --- Unary operation +SYCL_EXTERNAL auto test_unary_op(simd val) SYCL_ESIMD_FUNCTION { +// CHECK: define dso_local spir_func void @_Z13test_unary_op{{.*}}( +// CHECK: {{.*}} %[[RET_VEC_ADDR:[a-zA-Z0-9_\.]+]], +// CHECK: {{.*}} %[[VAL_PTR:[a-zA-Z0-9_\.]+]]){{.*}} { +// CHECK-LABEL: entry: + return -val; +// CHECK-NEXT: %[[VAL_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL_PTR]] +// CHECK-NEXT: %[[VAL_VEC:[a-zA-Z0-9_\.]+]] = load <8 x half>{{.*}} %[[VAL_VEC_ADDR]] +// CHECK-NEXT: %[[RES:[a-zA-Z0-9_\.]+]] = fneg <8 x half> %[[VAL_VEC]] +// CHECK-NEXT: store <8 x half>{{.*}}%[[RES]], {{.*}}%[[RET_VEC_ADDR]] +// CHECK-NEXT: ret void +// CHECK-LABEL: } +} + +// --- Binary operation on pair +SYCL_EXTERNAL auto test_binary_op1(simd val1, simd val2) SYCL_ESIMD_FUNCTION { +// CHECK: define dso_local spir_func void @_Z15test_binary_op1{{.*}}( +// CHECK: {{[^,]*}} %[[RET_VEC_ADDR:[a-zA-Z0-9_\.]+]], +// CHECK: {{[^,]*}} %[[VAL1_PTR:[a-zA-Z0-9_\.]+]], +// CHECK: {{.*}} %[[VAL2_PTR:[a-zA-Z0-9_\.]+]]){{.*}} { +// CHECK-LABEL: entry: + return val1 + val2; +// CHECK-NEXT: %[[VAL1_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL1_PTR]] +// CHECK-NEXT: %[[VAL1_VEC:[a-zA-Z0-9_\.]+]] = load <8 x half>{{.*}} %[[VAL1_VEC_ADDR]] +// CHECK-NEXT: %[[VAL2_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL2_PTR]] +// CHECK-NEXT: %[[VAL2_VEC:[a-zA-Z0-9_\.]+]] = load <8 x half>{{.*}} %[[VAL2_VEC_ADDR]] +// CHECK-NEXT: %[[RES:[a-zA-Z0-9_\.]+]] = fadd <8 x half> %[[VAL1_VEC]], %[[VAL2_VEC]] +// CHECK-NEXT: store <8 x half>{{.*}}%[[RES]], {{.*}}%[[RET_VEC_ADDR]] +// CHECK-NEXT: ret void +// CHECK-LABEL: } +} + +// --- Binary operation on pair +// The integer operand is expected to be converted to half type. +SYCL_EXTERNAL auto test_binary_op2(simd val1, simd val2) SYCL_ESIMD_FUNCTION { +// CHECK: define dso_local spir_func void @_Z15test_binary_op2{{[^\(]*}}( +// CHECK: <8 x half>{{[^,]*}}* %[[RET_VEC_ADDR:[a-zA-Z0-9_\.]+]], +// CHECK: <8 x half>* %[[VAL1_PTR:[a-zA-Z0-9_\.]+]], +// CHECK: <8 x i64>* %[[VAL2_PTR:[a-zA-Z0-9_\.]+]]){{.*}} { +// CHECK-LABEL: entry: + return val1 + val2; +// CHECK-NEXT: %[[VAL1_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL1_PTR]] +// CHECK-NEXT: %[[VAL1_VEC:[a-zA-Z0-9_\.]+]] = load <8 x half>{{.*}} %[[VAL1_VEC_ADDR]] +// CHECK-NEXT: %[[VAL2_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL2_PTR]] +// CHECK-NEXT: %[[VAL2_VEC:[a-zA-Z0-9_\.]+]] = load <8 x i64>{{.*}} %[[VAL2_VEC_ADDR]] +// CHECK-NEXT: %[[CONV:[a-zA-Z0-9_\.]+]] = sitofp <8 x i64> %[[VAL2_VEC]] to <8 x half> +// CHECK-NEXT: %[[RES:[a-zA-Z0-9_\.]+]] = fadd <8 x half> %[[VAL1_VEC]], %[[CONV]] +// CHECK-NEXT: store <8 x half>{{.*}}%[[RES]], {{.*}}%[[RET_VEC_ADDR]] +// CHECK-NEXT: ret void +// CHECK-LABEL: } +} + +// --- Comparison operation on pair +// The integer operand is expected to be converted to half type. +SYCL_EXTERNAL auto test_cmp_op(simd val1, simd val2) SYCL_ESIMD_FUNCTION { +// CHECK: define dso_local spir_func void @_Z11test_cmp_op{{[^\(]*}}( +// CHECK: <8 x i16>{{[^,]*}}* %[[RET_VEC_ADDR:[a-zA-Z0-9_\.]+]], +// CHECK: <8 x half>* %[[VAL1_PTR:[a-zA-Z0-9_\.]+]], +// CHECK: <8 x i64>* %[[VAL2_PTR:[a-zA-Z0-9_\.]+]]){{.*}} { +// CHECK-LABEL: entry: + return val1 < val2; +// CHECK-NEXT: %[[VAL1_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL1_PTR]] +// CHECK-NEXT: %[[VAL1_VEC:[a-zA-Z0-9_\.]+]] = load <8 x half>{{.*}} %[[VAL1_VEC_ADDR]] +// CHECK-NEXT: %[[VAL2_VEC_ADDR:[a-zA-Z0-9_\.]+]] = addrspacecast {{.*}} %[[VAL2_PTR]] +// CHECK-NEXT: %[[VAL2_VEC:[a-zA-Z0-9_\.]+]] = load <8 x i64>{{.*}} %[[VAL2_VEC_ADDR]] +// CHECK-NEXT: %[[CONV:[a-zA-Z0-9_\.]+]] = sitofp <8 x i64> %[[VAL2_VEC]] to <8 x half> +// CHECK-NEXT: %[[CMP:[a-zA-Z0-9_\.]+]] = fcmp olt <8 x half> %[[VAL1_VEC]], %[[CONV]] +// CHECK-NEXT: %[[RES:[a-zA-Z0-9_\.]+]] = zext <8 x i1> %[[CMP]] to <8 x i16> +// CHECK-NEXT: store <8 x i16>{{.*}}%[[RES]], {{.*}}%[[RET_VEC_ADDR]] +// CHECK-NEXT: ret void +// CHECK-LABEL: } +} From dc6ab69d5eaa8e59d341e932a97f3722fdc8ccf0 Mon Sep 17 00:00:00 2001 From: Konstantin S Bobrovsky Date: Sat, 1 Jan 2022 20:11:03 -0800 Subject: [PATCH 2/2] Fix clang-format Signed-off-by: Konstantin S Bobrovsky --- sycl/test/esimd/sycl_half_basic_ops.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/test/esimd/sycl_half_basic_ops.cpp b/sycl/test/esimd/sycl_half_basic_ops.cpp index c74c5e4f23be9..1e589fa0e0c90 100644 --- a/sycl/test/esimd/sycl_half_basic_ops.cpp +++ b/sycl/test/esimd/sycl_half_basic_ops.cpp @@ -12,6 +12,7 @@ using namespace sycl::ext::intel::experimental::esimd; using namespace sycl::ext::intel::experimental; using namespace sycl; +// clang-format off // --- Unary operation SYCL_EXTERNAL auto test_unary_op(simd val) SYCL_ESIMD_FUNCTION { // CHECK: define dso_local spir_func void @_Z13test_unary_op{{.*}}( @@ -85,3 +86,4 @@ SYCL_EXTERNAL auto test_cmp_op(simd val1, simd val2 // CHECK-NEXT: ret void // CHECK-LABEL: } } +// clang-format on