Skip to content

Commit 8445dc1

Browse files
AlexeySachkovAlexeySachkov
authored
[SYCL] Refactor LIT tests for diagnostics (#3956)
Switched ESIMD diagnostic tests to use FileCheck instead of -verify. Added %fsycl-host-only substitution, which contains a minimal amount of compiler arguments needed to only perform SYCL source compilation in host mode. Updated SYCL LIT tests checking diagnostics for host APIs to use that new substitution. This patch is motivated by a few of things: 1. enabling integration footer unconditionally in [SYCL] Move get_spec_constant_symbolic_ID to a separate header #3777 2. the fact that -verify is intended to be used with clang -cc1 3. the fact that VerifyDiagnosticConsumer doesn't handle (or maybe not even designed to) line markers in input file All those things lead to LIT failures, because -verify is not able to properly read diagnostic messages emitted by FE invocation for host part of DPC++/SYCL programs and reports that some diagnostics are either "seen, but not expected" or "expected, but not seen". Note: diagnostic messages themself look just fine and correct when observed through clang -fsycl, clang -cc1 -fsycl-device-only and clang -cc1 -fsycl-is-host invocations.
1 parent 619c462 commit 8445dc1

14 files changed

+116
-89
lines changed

‎sycl/test/basic_tests/image_accessor_types.cpp

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
// RUN: not %clangxx -fsyntax-only -std=c++17 %s -I %sycl_include/sycl 2>&1 | FileCheck %s
1+
// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
2+
// RUN: %clangxx -fsyntax-only -fsycl -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
23
#include <CL/sycl.hpp>
34
#include <iostream>
45

@@ -12,12 +13,10 @@ int main() {
1213
myQueue.submit([&](handler &cgh) {
1314
accessor<float, 2, access::mode::read, access::target::image> NotValidType1(
1415
srcImage, cgh);
15-
// CHECK: The data type of an image accessor must be only cl_int4, cl_uint4,
16-
// CHECK-SAME: cl_float4 or cl_half4
16+
// expected-error@CL/sycl/accessor.hpp:* {{The data type of an image accessor must be only cl_int4, cl_uint4, cl_float4 or cl_half4}}
1717
accessor<int2, 2, access::mode::read, access::target::image> NotValidType2(
1818
srcImage, cgh);
19-
// CHECK: The data type of an image accessor must be only cl_int4, cl_uint4,
20-
// CHECK-SAME: cl_float4 or cl_half4
19+
// expected-error@CL/sycl/accessor.hpp:* {{The data type of an image accessor must be only cl_int4, cl_uint4, cl_float4 or cl_half4}}
2120
accessor<float4, 2, access::mode::read, access::target::image>
2221
ValidSYCLFloat(srcImage, cgh);
2322
accessor<int4, 2, access::mode::read, access::target::image> ValidSYCLInt(

‎sycl/test/basic_tests/implicit_conversion_error.cpp

Lines changed: 2 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,5 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
2-
//=- implicit_conversion_error.cpp - Unintended implicit conversion check -=//
3-
//
4-
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5-
// See https://llvm.org/LICENSE.txt for license information.
6-
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7-
//
8-
//===-----------------------------------------------------------------------===//
1+
// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
2+
93
#include <CL/sycl.hpp>
104

115
int main() {

‎sycl/test/basic_tests/range_error.cpp

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,4 @@
1-
// RUN: %clangxx -fsycl -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only
2-
//==--------------- range_error.cpp - SYCL range error test ----------------==//
3-
//
4-
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5-
// See https://llvm.org/LICENSE.txt for license information.
6-
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7-
//
8-
//===----------------------------------------------------------------------===//
1+
// RUN: %clangxx %s %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning
92
#include <CL/sycl.hpp>
103
#include <iostream>
114
#include <cassert>
@@ -33,6 +26,8 @@ int main() {
3326
assert(three_dim_range.get(2) ==2);
3427
assert(three_dim_range[2] ==2);
3528
cout << "three_dim_range passed " << endl;
36-
cl::sycl::range<1> one_dim_range_f1(64, 2, 4);//expected-error {{no matching constructor for initialization of 'cl::sycl::range<1>'}}
37-
cl::sycl::range<2> two_dim_range_f1(64);//expected-error {{no matching constructor for initialization of 'cl::sycl::range<2>'}}
29+
// expected-error@+1 {{no matching constructor for initialization of 'cl::sycl::range<1>'}}
30+
cl::sycl::range<1> one_dim_range_f1(64, 2, 4);
31+
// expected-error@+1 {{no matching constructor for initialization of 'cl::sycl::range<2>'}}
32+
cl::sycl::range<2> two_dim_range_f1(64);
3833
}

‎sycl/test/basic_tests/set_arg_error.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -Xclang -verify %s -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only
1+
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
22

33
#include <CL/sycl.hpp>
44

‎sycl/test/basic_tests/vectors/ctad_fail.cpp

Lines changed: 1 addition & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,4 @@
1-
// RUN: %clangxx -std=c++17 -fsyntax-only -Xclang -verify %s -I %sycl_include/sycl -Xclang -verify-ignore-unexpected
2-
//==--------------- ctad.cpp - SYCL vector CTAD fail test ------------------==//
3-
//
4-
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5-
// See https://llvm.org/LICENSE.txt for license information.
6-
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7-
//
8-
//===----------------------------------------------------------------------===//
1+
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected %s
92
#include <CL/sycl.hpp>
103

114
namespace sycl = cl::sycl;

‎sycl/test/esimd/enums.cpp

Lines changed: 24 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
// This test checks compilation of various ESIMD enum types. Those which are
44
// deprecated must produce deprecation messages.
@@ -8,21 +8,35 @@
88
using namespace sycl::ext::intel::experimental::esimd;
99

1010
void foo() SYCL_ESIMD_FUNCTION {
11-
// These should produce deprecation messages:
11+
// These should produce deprecation messages for both host and device
12+
// compilations:
1213
int x;
13-
// expected-warning@+2 {{deprecated}}
14-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
14+
// CHECK: enums.cpp:16{{.*}}warning: 'WAIT' is deprecated
15+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
1516
x = static_cast<int>(ESIMD_SBARRIER_WAIT);
16-
// expected-warning@+2 {{deprecated}}
17-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
17+
// CHECK: enums.cpp:19{{.*}}warning: 'ATOMIC_ADD' is deprecated
18+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
1819
x = static_cast<int>(EsimdAtomicOpType::ATOMIC_ADD);
19-
// expected-warning@+2 {{deprecated}}
20-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
20+
// CHECK: enums.cpp:22{{.*}}warning: 'ESIMD_R_ENABLE' is deprecated
21+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2122
x = static_cast<int>(ChannelMaskType::ESIMD_R_ENABLE);
22-
// expected-warning@+2 {{deprecated}}
23-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
23+
// CHECK: enums.cpp:25{{.*}}warning: 'GENX_NOSAT' is deprecated
24+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2425
x = static_cast<int>(GENX_NOSAT);
2526

27+
// A "border" between host and device compilations
28+
// CHECK-LABEL: 4 warnings generated
29+
30+
// For host:
31+
// CHECK: enums.cpp:16{{.*}}warning: 'WAIT' is deprecated
32+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
33+
// CHECK: enums.cpp:19{{.*}}warning: 'ATOMIC_ADD' is deprecated
34+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
35+
// CHECK: enums.cpp:22{{.*}}warning: 'ESIMD_R_ENABLE' is deprecated
36+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
37+
// CHECK: enums.cpp:25{{.*}}warning: 'GENX_NOSAT' is deprecated
38+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
39+
2640
// These should compile cleanly:
2741
x = static_cast<int>(split_barrier_action::wait);
2842
x = static_cast<int>(atomic_op::add);

‎sycl/test/esimd/flat_atomic.cpp

Lines changed: 16 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
// This test checks compilation of ESIMD atomic APIs. Those which are deprecated
44
// must produce deprecation messages.
@@ -15,8 +15,8 @@ void kernel0(accessor<uint32_t, 1, access::mode::read_write,
1515
access::target::global_buffer> &buf) SYCL_ESIMD_FUNCTION {
1616
simd<uint32_t, 32> offsets(0, 1);
1717

18-
// expected-warning@+2 {{deprecated}}
19-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
18+
// CHECK: flat_atomic.cpp:20{{.*}}warning: 'ATOMIC_INC' is deprecated
19+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2020
flat_atomic<EsimdAtomicOpType::ATOMIC_INC, uint32_t, 32>(buf.get_pointer(), offsets, 1);
2121
flat_atomic<atomic_op::inc, uint32_t, 32>(buf.get_pointer(), offsets, 1);
2222
}
@@ -26,8 +26,8 @@ void kernel1(accessor<uint32_t, 1, access::mode::read_write,
2626
simd<uint32_t, 32> offsets(0, 1);
2727
simd<uint32_t, 32> v1(0, 1);
2828

29-
// expected-warning@+2 {{deprecated}}
30-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
29+
// CHECK: flat_atomic.cpp:31{{.*}}warning: 'ATOMIC_ADD' is deprecated
30+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
3131
flat_atomic<EsimdAtomicOpType::ATOMIC_ADD, uint32_t, 32>(buf.get_pointer(), offsets, v1, 1);
3232
flat_atomic<atomic_op::add, uint32_t, 32>(buf.get_pointer(), offsets, v1, 1);
3333
}
@@ -37,9 +37,18 @@ void kernel2(accessor<uint32_t, 1, access::mode::read_write,
3737
simd<uint32_t, 32> offsets(0, 1);
3838
simd<uint32_t, 32> v1(0, 1);
3939

40-
// expected-warning@+2 {{deprecated}}
41-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
40+
// CHECK: flat_atomic.cpp:42{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
41+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
4242
flat_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(buf.get_pointer(), offsets, v1, v1, 1);
4343
flat_atomic<atomic_op::cmpxchg, uint32_t, 32>(buf.get_pointer(), offsets, v1,
4444
v1, 1);
4545
}
46+
47+
// A "border" between device and host compilations
48+
// CHECK-LABEL: 3 warnings generated
49+
// CHECK: flat_atomic.cpp:20{{.*}}warning: 'ATOMIC_INC' is deprecated
50+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
51+
// CHECK: flat_atomic.cpp:31{{.*}}warning: 'ATOMIC_ADD' is deprecated
52+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
53+
// CHECK: flat_atomic.cpp:42{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
54+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:

‎sycl/test/esimd/gather4_scatter4.cpp

Lines changed: 20 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
// This test checks compilation of ESIMD slm gather4/scatter4 APIs. Those which
44
// are deprecated must produce deprecation messages.
@@ -16,23 +16,34 @@ void kernel(accessor<int, 1, access::mode::read_write,
1616
simd<uint32_t, 32> offsets(0, 1);
1717
simd<int, 32 * 4> v1(0, 1);
1818

19-
// expected-warning@+2 {{deprecated}}
20-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
19+
// CHECK: gather4_scatter4.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
20+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2121
auto v0 = gather4<int, 32, ESIMD_ABGR_ENABLE>(buf.get_pointer(), offsets);
22-
// expected-warning@+2 {{deprecated}}
23-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
22+
// CHECK: gather4_scatter4.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
23+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2424
v0 = gather4<int, 32, ChannelMaskType::ESIMD_ABGR_ENABLE>(buf.get_pointer(),
2525
offsets);
2626
v0 = gather4<int, 32, rgba_channel_mask::ABGR>(buf.get_pointer(), offsets);
2727

2828
v0 = v0 + v1;
2929

30-
// expected-warning@+2 {{deprecated}}
31-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
30+
// CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
31+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
3232
scatter4<int, 32, ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0, offsets);
33-
// expected-warning@+2 {{deprecated}}
34-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
33+
// CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
34+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
3535
scatter4<int, 32, ChannelMaskType::ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0,
3636
offsets);
3737
scatter4<int, 32, rgba_channel_mask::ABGR>(buf.get_pointer(), v0, offsets);
3838
}
39+
40+
// A "border" between host and device compilations
41+
// CHECK-LABEL: 4 warnings generated
42+
// CHECK: gather4_scatter4.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
43+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
44+
// CHECK: gather4_scatter4.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
45+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
46+
// CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
47+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
48+
// CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
49+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:

‎sycl/test/esimd/simd_copy_to_copy_from.cpp

Lines changed: 6 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
2+
// RUN: not %clangxx %fsycl-host-only -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
23

34
// This test checks that both host and device compilers can:
45
// - successfully compile simd::copy_to and simd::copy_from APIs
@@ -41,14 +42,10 @@ kernel3(accessor<int, 1, access::mode::read_write, access::target::local> &buf)
4142
SYCL_ESIMD_FUNCTION {
4243
simd<int, 32> v1(0, 1);
4344
simd<int, 32> v0;
44-
// expected-error@+3 {{no matching member function for call to 'copy_from'}}
45-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
46-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
45+
// CHECK: simd_copy_to_copy_from.cpp:46{{.*}}error: no matching member function for call to 'copy_from'
4746
v0.copy_from(buf, 0);
4847
v0 = v0 + v1;
49-
// expected-error@+3 {{no matching member function for call to 'copy_to'}}
50-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
51-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
48+
// CHECK: simd_copy_to_copy_from.cpp:49{{.*}}error: no matching member function for call to 'copy_to'
5249
v0.copy_to(buf, 0);
5350
}
5451

@@ -57,9 +54,7 @@ SYCL_EXTERNAL void kernel4(
5754
accessor<int, 1, access::mode::write, access::target::global_buffer> &buf)
5855
SYCL_ESIMD_FUNCTION {
5956
simd<int, 32> v;
60-
// expected-error@+3 {{no matching member function for call to 'copy_from'}}
61-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
62-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
57+
// CHECK: simd_copy_to_copy_from.cpp:58{{.*}}error: no matching member function for call to 'copy_from'
6358
v.copy_from(buf, 0);
6459
}
6560

@@ -68,8 +63,6 @@ SYCL_EXTERNAL void kernel5(
6863
accessor<int, 1, access::mode::read, access::target::global_buffer> &buf)
6964
SYCL_ESIMD_FUNCTION {
7065
simd<int, 32> v(0, 1);
71-
// expected-error@+3 {{no matching member function for call to 'copy_to'}}
72-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
73-
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
66+
// CHECK: simd_copy_to_copy_from.cpp:67{{.*}}error: no matching member function for call to 'copy_to'
7467
v.copy_to(buf, 0);
7568
}

‎sycl/test/esimd/simd_subscript.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: not %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
#include <sycl/ext/intel/experimental/esimd.hpp>
44

@@ -68,7 +68,8 @@ void test_simd_writable_subscript() SYCL_ESIMD_FUNCTION {
6868
void test_simd_const_subscript() SYCL_ESIMD_FUNCTION {
6969
const simd<int, 4> cv = 1;
7070
int val2 = cv[0]; // returns int instead of simd_view
71-
cv[1] = 0; // expected-error {{expression is not assignable}}
71+
// CHECK: simd_subscript.cpp:72{{.*}}error: expression is not assignable
72+
cv[1] = 0;
7273
}
7374

7475
void test_simd_view_assign_op() SYCL_ESIMD_FUNCTION {

‎sycl/test/esimd/slm_atomic.cpp

Lines changed: 16 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
// This test checks compilation of ESIMD slm atomic APIs. Those which are
44
// deprecated must produce deprecation messages.
@@ -14,8 +14,8 @@ using namespace cl::sycl;
1414
void kernel0() SYCL_ESIMD_FUNCTION {
1515
simd<uint32_t, 32> offsets(0, 1);
1616

17-
// expected-warning@+2 {{deprecated}}
18-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
17+
// CHECK: slm_atomic.cpp:19{{.*}}warning: 'ATOMIC_INC' is deprecated
18+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
1919
slm_atomic<EsimdAtomicOpType::ATOMIC_INC, uint32_t, 32>(offsets, 1);
2020
slm_atomic<atomic_op::inc, uint32_t, 32>(offsets, 1);
2121
}
@@ -24,8 +24,8 @@ void kernel1() SYCL_ESIMD_FUNCTION {
2424
simd<uint32_t, 32> offsets(0, 1);
2525
simd<uint32_t, 32> v1(0, 1);
2626

27-
// expected-warning@+2 {{deprecated}}
28-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
27+
// CHECK: slm_atomic.cpp:29{{.*}}warning: 'ATOMIC_ADD' is deprecated
28+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
2929
slm_atomic<EsimdAtomicOpType::ATOMIC_ADD, uint32_t, 32>(offsets, v1, 1);
3030
slm_atomic<atomic_op::add, uint32_t, 32>(offsets, v1, 1);
3131
}
@@ -34,8 +34,17 @@ void kernel2() SYCL_ESIMD_FUNCTION {
3434
simd<uint32_t, 32> offsets(0, 1);
3535
simd<uint32_t, 32> v1(0, 1);
3636

37-
// expected-warning@+2 {{deprecated}}
38-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
37+
// CHECK: slm_atomic.cpp:39{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
38+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
3939
slm_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(offsets, v1, v1, 1);
4040
slm_atomic<atomic_op::cmpxchg, uint32_t, 32>(offsets, v1, v1, 1);
4141
}
42+
43+
// A "border" between device and host compilations
44+
// CHECK-LABEL: 3 warnings generated
45+
// CHECK: slm_atomic.cpp:19{{.*}}warning: 'ATOMIC_INC' is deprecated
46+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
47+
// CHECK: slm_atomic.cpp:29{{.*}}warning: 'ATOMIC_ADD' is deprecated
48+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
49+
// CHECK: slm_atomic.cpp:39{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
50+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:

‎sycl/test/esimd/slm_load4.cpp

Lines changed: 12 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
1+
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
22

33
// This test checks compilation of ESIMD slm load4/store4 APIs. Those which are
44
// deprecated must produce deprecation messages.
@@ -14,15 +14,22 @@ void caller() SYCL_ESIMD_FUNCTION {
1414

1515
slm_init(1024);
1616

17-
// expected-warning@+2 {{deprecated}}
18-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
17+
// CHECK: slm_load4.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
18+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
1919
auto v0 = slm_load4<int, 32, ESIMD_ABGR_ENABLE>(offsets);
2020
v0 = slm_load4<int, 32, rgba_channel_mask::ABGR>(offsets);
2121

2222
v0 = v0 + v1;
2323

24-
// expected-warning@+2 {{deprecated}}
25-
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
24+
// CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
25+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
2626
slm_store4<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
2727
slm_store4<int, 32, rgba_channel_mask::ABGR>(v0, offsets);
2828
}
29+
30+
// A "border" between host and device compilations
31+
// CHECK-LABEL: 2 warnings generated
32+
// CHECK: slm_load4.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
33+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
34+
// CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
35+
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:

‎sycl/test/lit.cfg.py

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -90,6 +90,8 @@
9090
config.substitutions.append( ('%llvm_build_lib_dir', config.llvm_build_lib_dir ) )
9191
config.substitutions.append( ('%llvm_build_bin_dir', config.llvm_build_bin_dir ) )
9292

93+
config.substitutions.append( ('%fsycl-host-only', '-std=c++17 -Xclang -fsycl-is-host -isystem %s -isystem %s -isystem %s' % (config.sycl_include, config.opencl_include_dir, config.sycl_include + '/sycl/') ) )
94+
9395
llvm_config.add_tool_substitutions(['llvm-spirv'], [config.sycl_tools_dir])
9496

9597
config.substitutions.append( ('%RUN_ON_HOST', "env SYCL_DEVICE_FILTER=host ") )

0 commit comments

Comments
 (0)