[ESIMD] Add tests for ESIMDOptimizeVecArgCallConv pass. (#1289)

* [ESIMD] Add tests for ESIMDOptimizeVecArgCallConv pass.

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Author: kbobrovs

SYCL/ESIMD/vadd_usm_opqptr.cpp

@@ -0,0 +1,18 @@
+//==---------------- vadd_usm_opqptr.cpp  - DPC++ ESIMD on-device test -----==//
+//
+// 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: gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -Xclang -opaque-pointers %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+// TODO Running existing tests in opaque pointer mode should be supported by the
+// CI.
+
+// The test checks if vadd_usm.cpp works in opaque pointer mode.
+
+#include "vadd_usm.cpp"

SYCL/ESIMD/vec_arg_call_conv_ext.cpp

@@ -0,0 +1,87 @@
+//==-------------- vec_arg_call_conv_ext.cpp  - DPC++ ESIMD feature test ---==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Intel GPU is not really required, but the test does not make sense for
+// others.
+// REQUIRES: gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
+// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
+// RUN: FileCheck --input-file=%t.out.ll %s
+
+// Checks that ESIMDOptimizeVecArgCallConv does the right job as
+// a part of sycl-post-link.
+
+#include <sycl/ext/intel/esimd.hpp>
+
+using namespace sycl::ext::intel::esimd;
+
+// clang-format off
+
+//------------------------
+// Test1: Optimized parameter interleaves non - optimizeable ones.
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<int, 8> callee__sret__x_param_x(int i, simd<int, 8> x, int j) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <8 x i32> @_Z23callee__sret__x_param_x{{.*}}(i32 noundef %{{.*}}, <8 x i32> %{{.*}}, i32 noundef %{{.*}})
+  return x + (i + j);
+}
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<int, 8> test__sret__x_param_x(simd<int, 8> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <8 x i32> @_Z21test__sret__x_param_x{{.*}}(<8 x i32> %{{.*}})
+  return callee__sret__x_param_x(2, x, 1);
+// CHECK:  %{{.*}} = call spir_func <8 x i32> @_Z23callee__sret__x_param_x{{.*}}(i32 2, <8 x i32> %{{.*}}, i32 1)
+}
+
+//------------------------
+// Test2: "2-level fall through"
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<double, 32> callee__all_fall_through0(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <32 x double> @_Z25callee__all_fall_through0{{.*}}(<32 x double> %{{.*}})
+  return x;
+}
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<double, 32> callee__all_fall_through1(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <32 x double> @_Z25callee__all_fall_through1{{.*}}(<32 x double> %{{.*}})
+  return callee__all_fall_through0(x);
+// CHECK:  %{{.*}} = call spir_func <32 x double> @_Z25callee__all_fall_through0{{.*}}(<32 x double> %{{.*}})
+}
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<double, 32> test__all_fall_through(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <32 x double> @_Z22test__all_fall_through{{.*}}(<32 x double> %{{.*}})
+  return callee__all_fall_through1(x);
+// CHECK:  %{{.*}} = call spir_func <32 x double> @_Z25callee__all_fall_through1{{.*}}(<32 x double> %{{.*}})
+}
+
+//------------------------
+// Test3. First argument is passed by reference and updated in the callee,
+// must not be optimized.
+
+__attribute__((noinline))
+SYCL_EXTERNAL void callee_void__noopt_opt(simd<int, 8>& x, simd<int, 8> y) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func void @_Z22callee_void__noopt_opt{{.*}}(ptr {{.*}} %{{.*}}, <8 x i32> %{{.*}})
+  x = x + y;
+}
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<int, 8> test__sret__noopt_opt(simd<int, 8> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <8 x i32> @_Z21test__sret__noopt_opt{{.*}}(ptr noundef %{{.*}})
+  callee_void__noopt_opt(x, x);
+// CHECK:  call spir_func void @_Z22callee_void__noopt_opt{{.*}}(ptr addrspace(4) %{{.*}}, <8 x i32> %{{.*}})
+  return x;
+}
+
+//------------------------
+
+// Dummy main to satisfy linker.
+int main() {
+  return 0;
+}

SYCL/ESIMD/vec_arg_call_conv_smoke.cpp

@@ -0,0 +1,56 @@
+//==------------ vec_arg_call_conv_smoke.cpp  - DPC++ ESIMD feature test ---==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// Intel GPU is not really required, but the test does not make sense for
+// others.
+// REQUIRES: gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
+// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
+// RUN: FileCheck --input-file=%t.out.ll %s
+
+// Performs a basic check that ESIMDOptimizeVecArgCallConv does the right job as
+// a part of sycl-post-link.
+
+#include <sycl/ext/intel/esimd.hpp>
+
+using namespace sycl::ext::intel::esimd;
+
+ESIMD_PRIVATE simd<float, 3 * 32 * 4> GRF;
+#define V(x, w, i) (x).template select<w, 1>(i)
+
+// clang-format off
+
+// "Fall-through case", incoming optimizeable parameter is just returned
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<float, 16> callee__sret__param(simd<float, 16> x) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %[[PARAM:.+]])
+  return x;
+}
+
+// * Caller 1: simd object is read from array
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<float, 16> test__sret__fall_through__arr(simd<float, 16> *x, int i) SYCL_ESIMD_FUNCTION {
+// CHECK: define dso_local spir_func <16 x float> @_Z29test__sret__fall_through__arr{{.*}}(ptr addrspace(4) noundef %[[PARAM0:.+]], i32 noundef %{{.*}})
+  return callee__sret__param(x[i]);
+// CHECK: %{{.*}} = call spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %{{.*}})
+}
+
+// * Caller 2 : simd object is read from a global
+
+__attribute__((noinline))
+SYCL_EXTERNAL simd<float, 16> test__sret__fall_through__glob() SYCL_ESIMD_FUNCTION {
+  return callee__sret__param(V(GRF, 16, 0));
+// CHECK: %{{.*}} = call spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %{{.*}})
+}
+
+// Dummy main to satisfy linker.
+int main() {
+  return 0;
+}