diff --git a/sycl/test/esimd/vec_arg_call_conv_ext.cpp b/sycl/test/esimd/vec_arg_call_conv_ext.cpp new file mode 100644 index 0000000000000..a4cdc137b7192 --- /dev/null +++ b/sycl/test/esimd/vec_arg_call_conv_ext.cpp @@ -0,0 +1,71 @@ +// 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 + +using namespace sycl::ext::intel::esimd; + +// clang-format off + +//------------------------ +// Test1: Optimized parameter interleaves non - optimizeable ones. + +__attribute__((noinline)) +SYCL_EXTERNAL simd callee__sret__x_param_x(int i, simd 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 test__sret__x_param_x(simd 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 callee__all_fall_through0(simd 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 callee__all_fall_through1(simd 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 test__all_fall_through(simd 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& x, simd 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 test__sret__noopt_opt(simd 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; +} + +//------------------------ diff --git a/sycl/test/esimd/vec_arg_call_conv_smoke.cpp b/sycl/test/esimd/vec_arg_call_conv_smoke.cpp new file mode 100644 index 0000000000000..8f1f184842a4b --- /dev/null +++ b/sycl/test/esimd/vec_arg_call_conv_smoke.cpp @@ -0,0 +1,40 @@ +// 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 + +using namespace sycl::ext::intel::esimd; + +ESIMD_PRIVATE simd GRF; +#define V(x, w, i) (x).template select(i) + +// clang-format off + +// "Fall-through case", incoming optimizeable parameter is just returned + +__attribute__((noinline)) +SYCL_EXTERNAL simd callee__sret__param(simd 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 test__sret__fall_through__arr(simd *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 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> %{{.*}}) +}