From b2f406db6e915beb96c1fcf9e9ee4ebb883b80f7 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 16 Nov 2022 08:34:43 -0500 Subject: [PATCH 1/2] [SYCL][NFC] Move ESIMD tests in-tree Tests were added to intel/llvm-test-suite repo (intel/llvm-test-suite#1289), but they are too low-level to be located in there. Moved them in-tree. --- sycl/test/esimd/vec_arg_call_conv_ext.cpp | 71 +++++++++++++++++++++ sycl/test/esimd/vec_arg_call_conv_smoke.cpp | 40 ++++++++++++ 2 files changed, 111 insertions(+) create mode 100644 sycl/test/esimd/vec_arg_call_conv_ext.cpp create mode 100644 sycl/test/esimd/vec_arg_call_conv_smoke.cpp 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..92acdf57375b3 --- /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 -c -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..f644e1a40d5c6 --- /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 -c -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> %{{.*}}) +} From 526af47f3545691e3a3af707ba318ae2a868cc11 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Thu, 17 Nov 2022 13:51:23 +0100 Subject: [PATCH 2/2] Apply suggestions from code review --- sycl/test/esimd/vec_arg_call_conv_ext.cpp | 2 +- sycl/test/esimd/vec_arg_call_conv_smoke.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/esimd/vec_arg_call_conv_ext.cpp b/sycl/test/esimd/vec_arg_call_conv_ext.cpp index 92acdf57375b3..a4cdc137b7192 100644 --- a/sycl/test/esimd/vec_arg_call_conv_ext.cpp +++ b/sycl/test/esimd/vec_arg_call_conv_ext.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -c -Xclang -emit-llvm -o %t.comp.ll %s +// 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 diff --git a/sycl/test/esimd/vec_arg_call_conv_smoke.cpp b/sycl/test/esimd/vec_arg_call_conv_smoke.cpp index f644e1a40d5c6..8f1f184842a4b 100644 --- a/sycl/test/esimd/vec_arg_call_conv_smoke.cpp +++ b/sycl/test/esimd/vec_arg_call_conv_smoke.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -c -Xclang -emit-llvm -o %t.comp.ll %s +// 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