Skip to content

[SYCL] Refactor LIT tests for diagnostics #3956

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
Jun 25, 2021
Merged

[SYCL] Refactor LIT tests for diagnostics #3956

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
0f8f481
[SYCL] Refactor LIT tests for diagnostics
AlexeySachkov Jun 16, 2021
118ec38
Apply review comments
AlexeySachkov Jun 21, 2021
0da8cdc
Apply comments
AlexeySachkov Jun 23, 2021
18158a4
temp
AlexeySachkov Jun 24, 2021
35b6ef9
Ensure that no unexpected diagnostics emitted with FileCheck
AlexeySachkov Jun 24, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 4 additions & 5 deletions sycl/test/basic_tests/image_accessor_types.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
// RUN: not %clangxx -fsyntax-only -std=c++17 %s -I %sycl_include/sycl 2>&1 | FileCheck %s
// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
// RUN: %clangxx -fsyntax-only -fsycl -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
#include <CL/sycl.hpp>
#include <iostream>

Expand All @@ -12,12 +13,10 @@ int main() {
myQueue.submit([&](handler &cgh) {
accessor<float, 2, access::mode::read, access::target::image> NotValidType1(
srcImage, cgh);
// CHECK: The data type of an image accessor must be only cl_int4, cl_uint4,
// CHECK-SAME: cl_float4 or cl_half4
// 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}}
accessor<int2, 2, access::mode::read, access::target::image> NotValidType2(
srcImage, cgh);
// CHECK: The data type of an image accessor must be only cl_int4, cl_uint4,
// CHECK-SAME: cl_float4 or cl_half4
// 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}}
accessor<float4, 2, access::mode::read, access::target::image>
ValidSYCLFloat(srcImage, cgh);
accessor<int4, 2, access::mode::read, access::target::image> ValidSYCLInt(
Expand Down
10 changes: 2 additions & 8 deletions sycl/test/basic_tests/implicit_conversion_error.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,5 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning
//=- implicit_conversion_error.cpp - Unintended implicit conversion check -=//
//
// 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
//
//===-----------------------------------------------------------------------===//
// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s

#include <CL/sycl.hpp>

int main() {
Expand Down
15 changes: 5 additions & 10 deletions sycl/test/basic_tests/range_error.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,4 @@
// RUN: %clangxx -fsycl -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only
//==--------------- range_error.cpp - SYCL range error 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
//
//===----------------------------------------------------------------------===//
// RUN: %clangxx %s %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning
#include <CL/sycl.hpp>
#include <iostream>
#include <cassert>
Expand Down Expand Up @@ -33,6 +26,8 @@ int main() {
assert(three_dim_range.get(2) ==2);
assert(three_dim_range[2] ==2);
cout << "three_dim_range passed " << endl;
cl::sycl::range<1> one_dim_range_f1(64, 2, 4);//expected-error {{no matching constructor for initialization of 'cl::sycl::range<1>'}}
cl::sycl::range<2> two_dim_range_f1(64);//expected-error {{no matching constructor for initialization of 'cl::sycl::range<2>'}}
// expected-error@+1 {{no matching constructor for initialization of 'cl::sycl::range<1>'}}
cl::sycl::range<1> one_dim_range_f1(64, 2, 4);
// expected-error@+1 {{no matching constructor for initialization of 'cl::sycl::range<2>'}}
cl::sycl::range<2> two_dim_range_f1(64);
}
2 changes: 1 addition & 1 deletion sycl/test/basic_tests/set_arg_error.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -Xclang -verify %s -I %sycl_include -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s

#include <CL/sycl.hpp>

Expand Down
9 changes: 1 addition & 8 deletions sycl/test/basic_tests/vectors/ctad_fail.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,11 +1,4 @@
// RUN: %clangxx -std=c++17 -fsyntax-only -Xclang -verify %s -I %sycl_include/sycl -Xclang -verify-ignore-unexpected
//==--------------- ctad.cpp - SYCL vector CTAD fail 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
//
//===----------------------------------------------------------------------===//
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected %s
#include <CL/sycl.hpp>

namespace sycl = cl::sycl;
Expand Down
34 changes: 24 additions & 10 deletions sycl/test/esimd/enums.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

void foo() SYCL_ESIMD_FUNCTION {
// These should produce deprecation messages:
// These should produce deprecation messages for both host and device
// compilations:
int x;
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: enums.cpp:16{{.*}}warning: 'WAIT' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
x = static_cast<int>(ESIMD_SBARRIER_WAIT);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: enums.cpp:19{{.*}}warning: 'ATOMIC_ADD' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
x = static_cast<int>(EsimdAtomicOpType::ATOMIC_ADD);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: enums.cpp:22{{.*}}warning: 'ESIMD_R_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
x = static_cast<int>(ChannelMaskType::ESIMD_R_ENABLE);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: enums.cpp:25{{.*}}warning: 'GENX_NOSAT' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
x = static_cast<int>(GENX_NOSAT);

// A "border" between host and device compilations
// CHECK-LABEL: 4 warnings generated

// For host:
// CHECK: enums.cpp:16{{.*}}warning: 'WAIT' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: enums.cpp:19{{.*}}warning: 'ATOMIC_ADD' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: enums.cpp:22{{.*}}warning: 'ESIMD_R_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: enums.cpp:25{{.*}}warning: 'GENX_NOSAT' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:

// These should compile cleanly:
x = static_cast<int>(split_barrier_action::wait);
x = static_cast<int>(atomic_op::add);
Expand Down
23 changes: 16 additions & 7 deletions sycl/test/esimd/flat_atomic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

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

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

// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: flat_atomic.cpp:42{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
flat_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(buf.get_pointer(), offsets, v1, v1, 1);
flat_atomic<atomic_op::cmpxchg, uint32_t, 32>(buf.get_pointer(), offsets, v1,
v1, 1);
}

// A "border" between device and host compilations
// CHECK-LABEL: 3 warnings generated
// CHECK: flat_atomic.cpp:20{{.*}}warning: 'ATOMIC_INC' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: flat_atomic.cpp:31{{.*}}warning: 'ATOMIC_ADD' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: flat_atomic.cpp:42{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
29 changes: 20 additions & 9 deletions sycl/test/esimd/gather4_scatter4.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

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

v0 = v0 + v1;

// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
scatter4<int, 32, ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0, offsets);
// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
scatter4<int, 32, ChannelMaskType::ESIMD_ABGR_ENABLE>(buf.get_pointer(), v0,
offsets);
scatter4<int, 32, rgba_channel_mask::ABGR>(buf.get_pointer(), v0, offsets);
}

// A "border" between host and device compilations
// CHECK-LABEL: 4 warnings generated
// CHECK: gather4_scatter4.cpp:21{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: gather4_scatter4.cpp:24{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: gather4_scatter4.cpp:32{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: gather4_scatter4.cpp:35{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
19 changes: 6 additions & 13 deletions sycl/test/esimd/simd_copy_to_copy_from.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"
// RUN: not %clangxx %fsycl-host-only -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

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

Expand All @@ -68,8 +63,6 @@ SYCL_EXTERNAL void kernel5(
accessor<int, 1, access::mode::read, access::target::global_buffer> &buf)
SYCL_ESIMD_FUNCTION {
simd<int, 32> v(0, 1);
// expected-error@+3 {{no matching member function for call to 'copy_to'}}
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
// expected-note@sycl/ext/intel/experimental/esimd/simd.hpp:* {{}}
// CHECK: simd_copy_to_copy_from.cpp:67{{.*}}error: no matching member function for call to 'copy_to'
v.copy_to(buf, 0);
}
5 changes: 3 additions & 2 deletions sycl/test/esimd/simd_subscript.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: not %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

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

void test_simd_view_assign_op() SYCL_ESIMD_FUNCTION {
Expand Down
23 changes: 16 additions & 7 deletions sycl/test/esimd/slm_atomic.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

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

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

// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: slm_atomic.cpp:39{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
slm_atomic<EsimdAtomicOpType::ATOMIC_CMPXCHG, uint32_t, 32>(offsets, v1, v1, 1);
slm_atomic<atomic_op::cmpxchg, uint32_t, 32>(offsets, v1, v1, 1);
}

// A "border" between device and host compilations
// CHECK-LABEL: 3 warnings generated
// CHECK: slm_atomic.cpp:19{{.*}}warning: 'ATOMIC_INC' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: slm_atomic.cpp:29{{.*}}warning: 'ATOMIC_ADD' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
// CHECK: slm_atomic.cpp:39{{.*}}warning: 'ATOMIC_CMPXCHG' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp{{.*}}note:
17 changes: 12 additions & 5 deletions sycl/test/esimd/slm_load4.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s
// RUN: %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s --implicit-check-not="warning:" --implicit-check-not="error:"

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

slm_init(1024);

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

v0 = v0 + v1;

// expected-warning@+2 {{deprecated}}
// expected-note@sycl/ext/intel/experimental/esimd/common.hpp:* {{}}
// CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
slm_store4<int, 32, ESIMD_ABGR_ENABLE>(v0, offsets);
slm_store4<int, 32, rgba_channel_mask::ABGR>(v0, offsets);
}

// A "border" between host and device compilations
// CHECK-LABEL: 2 warnings generated
// CHECK: slm_load4.cpp:19{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
// CHECK: slm_load4.cpp:26{{.*}}warning: 'ESIMD_ABGR_ENABLE' is deprecated
// CHECK: sycl/ext/intel/experimental/esimd/common.hpp:{{.*}}note:
2 changes: 2 additions & 0 deletions sycl/test/lit.cfg.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,6 +90,8 @@
config.substitutions.append( ('%llvm_build_lib_dir', config.llvm_build_lib_dir ) )
config.substitutions.append( ('%llvm_build_bin_dir', config.llvm_build_bin_dir ) )

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/') ) )
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not sure -std=c++17 needs to be part of fsycl-host-only, but that is probably Ok, if that is intentional and has meaning "using sycl switches assumes using c++17"

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

if that is intentional and has meaning "using sycl switches assumes using c++17"

That was intentional, we can't compile our headers in c++14 mode, unfortunately. It is probably a bug, but for now I think that it is ok to stick with c++17 explicitly


llvm_config.add_tool_substitutions(['llvm-spirv'], [config.sycl_tools_dir])

config.substitutions.append( ('%RUN_ON_HOST', "env SYCL_DEVICE_FILTER=host ") )
Expand Down
Loading