Skip to content

Commit bbd267a

Browse files
kbobrovskbobrovs
authored
Merge pull request intel#1247 from kbobrovs/scsew-cfl-04-testroot-topic
[ESIMD] Add tests for ESIMDOptimizeVecArgCallConv and opaque pointer support.
2 parents 8cc0a66 + 68aea0c commit bbd267a

12 files changed

+188
-0
lines changed

‎SYCL/ESIMD/vadd_usm_opqptr.cpp

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,18 @@
1+
//==---------------- vadd_usm_opqptr.cpp - DPC++ ESIMD on-device test -----==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// REQUIRES: gpu
9+
// UNSUPPORTED: cuda || hip
10+
// RUN: %clangxx -fsycl -Xclang -opaque-pointers %s -o %t.out
11+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
12+
13+
// TODO Running existing tests in opaque pointer mode should be supported by the
14+
// CI.
15+
16+
// The test checks if vadd_usm.cpp works in opaque pointer mode.
17+
18+
#include "vadd_usm.cpp"

‎SYCL/ESIMD/vec_arg_call_conv_ext.cpp

Lines changed: 87 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,87 @@
1+
//==-------------- vec_arg_call_conv_ext.cpp - DPC++ ESIMD feature test ---==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// Intel GPU is not really required, but the test does not make sense for
9+
// others.
10+
// REQUIRES: gpu
11+
// UNSUPPORTED: cuda || hip
12+
// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
13+
// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
14+
// RUN: FileCheck --input-file=%t.out.ll %s
15+
16+
// Checks that ESIMDOptimizeVecArgCallConv does the right job as
17+
// a part of sycl-post-link.
18+
19+
#include <sycl/ext/intel/esimd.hpp>
20+
21+
using namespace sycl::ext::intel::esimd;
22+
23+
// clang-format off
24+
25+
//------------------------
26+
// Test1: Optimized parameter interleaves non - optimizeable ones.
27+
28+
__attribute__((noinline))
29+
SYCL_EXTERNAL simd<int, 8> callee__sret__x_param_x(int i, simd<int, 8> x, int j) SYCL_ESIMD_FUNCTION {
30+
// CHECK: define dso_local spir_func <8 x i32> @_Z23callee__sret__x_param_x{{.*}}(i32 noundef %{{.*}}, <8 x i32> %{{.*}}, i32 noundef %{{.*}})
31+
return x + (i + j);
32+
}
33+
34+
__attribute__((noinline))
35+
SYCL_EXTERNAL simd<int, 8> test__sret__x_param_x(simd<int, 8> x) SYCL_ESIMD_FUNCTION {
36+
// CHECK: define dso_local spir_func <8 x i32> @_Z21test__sret__x_param_x{{.*}}(<8 x i32> %{{.*}})
37+
return callee__sret__x_param_x(2, x, 1);
38+
// CHECK: %{{.*}} = call spir_func <8 x i32> @_Z23callee__sret__x_param_x{{.*}}(i32 2, <8 x i32> %{{.*}}, i32 1)
39+
}
40+
41+
//------------------------
42+
// Test2: "2-level fall through"
43+
44+
__attribute__((noinline))
45+
SYCL_EXTERNAL simd<double, 32> callee__all_fall_through0(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
46+
// CHECK: define dso_local spir_func <32 x double> @_Z25callee__all_fall_through0{{.*}}(<32 x double> %{{.*}})
47+
return x;
48+
}
49+
50+
__attribute__((noinline))
51+
SYCL_EXTERNAL simd<double, 32> callee__all_fall_through1(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
52+
// CHECK: define dso_local spir_func <32 x double> @_Z25callee__all_fall_through1{{.*}}(<32 x double> %{{.*}})
53+
return callee__all_fall_through0(x);
54+
// CHECK: %{{.*}} = call spir_func <32 x double> @_Z25callee__all_fall_through0{{.*}}(<32 x double> %{{.*}})
55+
}
56+
57+
__attribute__((noinline))
58+
SYCL_EXTERNAL simd<double, 32> test__all_fall_through(simd<double, 32> x) SYCL_ESIMD_FUNCTION {
59+
// CHECK: define dso_local spir_func <32 x double> @_Z22test__all_fall_through{{.*}}(<32 x double> %{{.*}})
60+
return callee__all_fall_through1(x);
61+
// CHECK: %{{.*}} = call spir_func <32 x double> @_Z25callee__all_fall_through1{{.*}}(<32 x double> %{{.*}})
62+
}
63+
64+
//------------------------
65+
// Test3. First argument is passed by reference and updated in the callee,
66+
// must not be optimized.
67+
68+
__attribute__((noinline))
69+
SYCL_EXTERNAL void callee_void__noopt_opt(simd<int, 8>& x, simd<int, 8> y) SYCL_ESIMD_FUNCTION {
70+
// CHECK: define dso_local spir_func void @_Z22callee_void__noopt_opt{{.*}}(ptr {{.*}} %{{.*}}, <8 x i32> %{{.*}})
71+
x = x + y;
72+
}
73+
74+
__attribute__((noinline))
75+
SYCL_EXTERNAL simd<int, 8> test__sret__noopt_opt(simd<int, 8> x) SYCL_ESIMD_FUNCTION {
76+
// CHECK: define dso_local spir_func <8 x i32> @_Z21test__sret__noopt_opt{{.*}}(ptr noundef %{{.*}})
77+
callee_void__noopt_opt(x, x);
78+
// CHECK: call spir_func void @_Z22callee_void__noopt_opt{{.*}}(ptr addrspace(4) %{{.*}}, <8 x i32> %{{.*}})
79+
return x;
80+
}
81+
82+
//------------------------
83+
84+
// Dummy main to satisfy linker.
85+
int main() {
86+
return 0;
87+
}

‎SYCL/ESIMD/vec_arg_call_conv_smoke.cpp

Lines changed: 56 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,56 @@
1+
//==------------ vec_arg_call_conv_smoke.cpp - DPC++ ESIMD feature test ---==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// Intel GPU is not really required, but the test does not make sense for
9+
// others.
10+
// REQUIRES: gpu
11+
// UNSUPPORTED: cuda || hip
12+
// RUN: %clangxx -fsycl -Xclang -opaque-pointers -fsycl-device-only -Xclang -emit-llvm -o %t.comp.ll %s
13+
// RUN: sycl-post-link -ir-output-only -lower-esimd -S %t.comp.ll -o %t.out.ll
14+
// RUN: FileCheck --input-file=%t.out.ll %s
15+
16+
// Performs a basic check that ESIMDOptimizeVecArgCallConv does the right job as
17+
// a part of sycl-post-link.
18+
19+
#include <sycl/ext/intel/esimd.hpp>
20+
21+
using namespace sycl::ext::intel::esimd;
22+
23+
ESIMD_PRIVATE simd<float, 3 * 32 * 4> GRF;
24+
#define V(x, w, i) (x).template select<w, 1>(i)
25+
26+
// clang-format off
27+
28+
// "Fall-through case", incoming optimizeable parameter is just returned
29+
30+
__attribute__((noinline))
31+
SYCL_EXTERNAL simd<float, 16> callee__sret__param(simd<float, 16> x) SYCL_ESIMD_FUNCTION {
32+
// CHECK: define dso_local spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %[[PARAM:.+]])
33+
return x;
34+
}
35+
36+
// * Caller 1: simd object is read from array
37+
38+
__attribute__((noinline))
39+
SYCL_EXTERNAL simd<float, 16> test__sret__fall_through__arr(simd<float, 16> *x, int i) SYCL_ESIMD_FUNCTION {
40+
// CHECK: define dso_local spir_func <16 x float> @_Z29test__sret__fall_through__arr{{.*}}(ptr addrspace(4) noundef %[[PARAM0:.+]], i32 noundef %{{.*}})
41+
return callee__sret__param(x[i]);
42+
// CHECK: %{{.*}} = call spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %{{.*}})
43+
}
44+
45+
// * Caller 2 : simd object is read from a global
46+
47+
__attribute__((noinline))
48+
SYCL_EXTERNAL simd<float, 16> test__sret__fall_through__glob() SYCL_ESIMD_FUNCTION {
49+
return callee__sret__param(V(GRF, 16, 0));
50+
// CHECK: %{{.*}} = call spir_func <16 x float> @_Z19callee__sret__param{{.*}}(<16 x float> %{{.*}})
51+
}
52+
53+
// Dummy main to satisfy linker.
54+
int main() {
55+
return 0;
56+
}

‎config_sycl/esimd_vadd_usm_opqptr.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/ESIMD/vadd_usm_opqptr.cpp

‎config_sycl/esimd_vadd_usm_opqptr.xml

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
<?xml version="1.0" encoding="UTF-8" ?>
2+
<test name="esimd_vadd_usm_opqptr" driverID="llvm_test_suite_sycl">
3+
<description>WARNING: DON'T UPDATE THIS FILE MANUALLY!!!
4+
This config file auto-generated by suite_generator_sycl.pl.</description>
5+
</test>

‎config_sycl/esimd_vec_arg_call_conv_ext.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/ESIMD/vec_arg_call_conv_ext.cpp

‎config_sycl/esimd_vec_arg_call_conv_ext.xml

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
<?xml version="1.0" encoding="UTF-8" ?>
2+
<test name="esimd_vec_arg_call_conv_ext" driverID="llvm_test_suite_sycl">
3+
<description>WARNING: DON'T UPDATE THIS FILE MANUALLY!!!
4+
This config file auto-generated by suite_generator_sycl.pl.</description>
5+
</test>

‎config_sycl/esimd_vec_arg_call_conv_smoke.info

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
SYCL/ESIMD/vec_arg_call_conv_smoke.cpp

‎config_sycl/esimd_vec_arg_call_conv_smoke.xml

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,5 @@
1+
<?xml version="1.0" encoding="UTF-8" ?>
2+
<test name="esimd_vec_arg_call_conv_smoke" driverID="llvm_test_suite_sycl">
3+
<description>WARNING: DON'T UPDATE THIS FILE MANUALLY!!!
4+
This config file auto-generated by suite_generator_sycl.pl.</description>
5+
</test>

‎llvm_test_suite_sycl.xml

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -563,6 +563,9 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
563563
<test configFile="config_sycl/esimd_vadd_half.xml" splitGroup="esimd" testName="esimd_vadd_half" />
564564
<test configFile="config_sycl/esimd_vadd_raw_send.xml" splitGroup="esimd" testName="esimd_vadd_raw_send" />
565565
<test configFile="config_sycl/esimd_vadd_usm.xml" splitGroup="esimd" testName="esimd_vadd_usm" />
566+
<test configFile="config_sycl/esimd_vadd_usm_opqptr.xml" splitGroup="esimd" testName="esimd_vadd_usm_opqptr" />
567+
<test configFile="config_sycl/esimd_vec_arg_call_conv_ext.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_ext" />
568+
<test configFile="config_sycl/esimd_vec_arg_call_conv_smoke.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_smoke" />
566569
<test configFile="config_sycl/filterselector_filter_list_cpu_gpu_acc.xml" splitGroup="filterselector" testName="filterselector_filter_list_cpu_gpu_acc" />
567570
<test configFile="config_sycl/filterselector_reuse.xml" splitGroup="filterselector" testName="filterselector_reuse" />
568571
<test configFile="config_sycl/filterselector_select.xml" splitGroup="filterselector" testName="filterselector_select" />

‎llvm_test_suite_sycl_esimd.xml

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -267,5 +267,8 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
267267
<test configFile="config_sycl/esimd_vadd_half.xml" splitGroup="esimd" testName="esimd_vadd_half" />
268268
<test configFile="config_sycl/esimd_vadd_raw_send.xml" splitGroup="esimd" testName="esimd_vadd_raw_send" />
269269
<test configFile="config_sycl/esimd_vadd_usm.xml" splitGroup="esimd" testName="esimd_vadd_usm" />
270+
<test configFile="config_sycl/esimd_vadd_usm_opqptr.xml" splitGroup="esimd" testName="esimd_vadd_usm_opqptr" />
271+
<test configFile="config_sycl/esimd_vec_arg_call_conv_ext.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_ext" />
272+
<test configFile="config_sycl/esimd_vec_arg_call_conv_smoke.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_smoke" />
270273
</tests>
271274
</suite>

‎llvm_test_suite_sycl_valgrind.xml

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -563,6 +563,9 @@ Sources repo https://github.com/intel-innersource/applications.compilers.tests.l
563563
<test configFile="config_sycl/esimd_vadd_half.xml" splitGroup="esimd" testName="esimd_vadd_half" />
564564
<test configFile="config_sycl/esimd_vadd_raw_send.xml" splitGroup="esimd" testName="esimd_vadd_raw_send" />
565565
<test configFile="config_sycl/esimd_vadd_usm.xml" splitGroup="esimd" testName="esimd_vadd_usm" />
566+
<test configFile="config_sycl/esimd_vadd_usm_opqptr.xml" splitGroup="esimd" testName="esimd_vadd_usm_opqptr" />
567+
<test configFile="config_sycl/esimd_vec_arg_call_conv_ext.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_ext" />
568+
<test configFile="config_sycl/esimd_vec_arg_call_conv_smoke.xml" splitGroup="esimd" testName="esimd_vec_arg_call_conv_smoke" />
566569
<test configFile="config_sycl/filterselector_filter_list_cpu_gpu_acc.xml" splitGroup="filterselector" testName="filterselector_filter_list_cpu_gpu_acc" />
567570
<test configFile="config_sycl/filterselector_reuse.xml" splitGroup="filterselector" testName="filterselector_reuse" />
568571
<test configFile="config_sycl/filterselector_select.xml" splitGroup="filterselector" testName="filterselector_select" />

0 commit comments

Comments
 (0)