Skip to content

Commit 7744f69

Browse files
jchlandajchlanda
committed
Remove SupportsNonconformingLambdaSyntax and adjust tests
1 parent 9431687 commit 7744f69

File tree

5 files changed

+53
-61
lines changed

5 files changed

+53
-61
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1589,7 +1589,6 @@ def SYCLIntelMinWorkGroupsPerComputeUnit : InheritableAttr {
15891589
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
15901590
let Subjects = SubjectList<[Function], ErrorDiag>;
15911591
let Documentation = [SYCLIntelMinWorkGroupsPerComputeUnitAttrDocs];
1592-
let SupportsNonconformingLambdaSyntax = 1;
15931592
}
15941593

15951594
def SYCLIntelMaxWorkGroupsPerMultiprocessor : InheritableAttr {
@@ -1598,7 +1597,6 @@ def SYCLIntelMaxWorkGroupsPerMultiprocessor : InheritableAttr {
15981597
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
15991598
let Subjects = SubjectList<[Function], ErrorDiag>;
16001599
let Documentation = [SYCLIntelMaxWorkGroupsPerMultiprocessorDocs];
1601-
let SupportsNonconformingLambdaSyntax = 1;
16021600
}
16031601

16041602
def SYCLIntelMaxGlobalWorkDim : InheritableAttr {

‎clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -42,9 +42,9 @@ int main() {
4242

4343
// Test attribute is applied on lambda.
4444
h.single_task<class kernel_name2>(
45-
[]() [[intel::max_work_group_size(8, 8, 8),
46-
intel::min_work_groups_per_cu(2),
47-
intel::max_work_groups_per_mp(4)]] {});
45+
[] [[intel::max_work_group_size(8, 8, 8),
46+
intel::min_work_groups_per_cu(2),
47+
intel::max_work_groups_per_mp(4)]] () {});
4848

4949
// Test class template argument.
5050
Functor<6> f;

‎clang/test/CodeGenSYCL/lb_sm_90.cpp

Lines changed: 7 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_90 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
2-
// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -verify
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_90 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK-IR
2+
// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -Wno-c++23-extensions -verify
33
// expected-no-diagnostics
44

55
// Maximum work groups per multi-processor, mapped to maxclusterrank PTX
@@ -21,13 +21,11 @@ int main() {
2121
sycl::range<1> Gws(32);
2222

2323
Q.submit([&](sycl::handler &cgh) {
24-
cgh.parallel_for<class K1>(Gws,
25-
[=](sycl::id<1>) [[intel::max_work_group_size(1, 1, 256),
26-
intel::min_work_groups_per_cu(2),
27-
intel::max_work_groups_per_mp(4)]] {
28-
volatile int A = 42;
29-
});
30-
});
24+
cgh.parallel_for<class K1>(Gws, [=] [[intel::max_work_group_size(1, 1, 256),
25+
intel::min_work_groups_per_cu(2),
26+
intel::max_work_groups_per_mp(4)]] (
27+
sycl::id<1>) { volatile int A = 42; });
28+
});
3129
// CHECK-IR: !min_work_groups_per_cu [[MWGPCU:![0-9]+]]
3230
// CHECK-IR: !max_work_groups_per_mp [[MWGPMP:![0-9]+]]
3331
// CHECK-IR: !max_work_group_size [[MWGS:![0-9]+]]

‎clang/test/SemaSYCL/lb_sm_70.cpp

Lines changed: 33 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -S -emit-llvm %s -o -ferror-limit=100 -fsyntax-only -verify %s
1+
// RUN: %clang_cc1 -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_70 -fsycl-is-device -Wno-c++23-extensions -S -emit-llvm %s -o -ferror-limit=100 -fsyntax-only -verify %s
22

33
// Maximum work groups per multi-processor, mapped to maxclusterrank PTX
44
// directive, is an SM_90 feature, make sure that correct warning is issued on
@@ -19,41 +19,38 @@ int main() {
1919
sycl::queue Q{};
2020

2121
Q.submit([&](sycl::handler &cgh) {
22-
// expected-warning@+5 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
23-
cgh.single_task<class T1>(
24-
[=]()
25-
[[intel::max_work_group_size(1, 1, 256),
26-
intel::min_work_groups_per_cu(2),
27-
intel::max_work_groups_per_mp(4)]] { volatile int A = 42; });
28-
29-
constexpr float A = 2.0;
30-
// expected-error@+5 {{'min_work_groups_per_cu' attribute requires an integer constant}}
31-
// expected-warning@+5 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
32-
cgh.single_task<class T2>(
33-
[=]()
34-
[[intel::max_work_group_size(1, 1, 256),
35-
intel::min_work_groups_per_cu(A),
36-
intel::max_work_groups_per_mp(4)]] { volatile int A = 42; });
37-
38-
// expected-error@+3 {{'min_work_groups_per_cu' attribute requires an integer constant}}
39-
cgh.single_task<class T3>(
40-
[=]() [[intel::max_work_group_size(1, 1, 256),
41-
intel::min_work_groups_per_cu(2147483647 + 1)]]
42-
{ volatile int A = 42; });
43-
44-
// expected-warning@+4 {{attribute 'min_work_groups_per_cu' is already applied with different arguments}}
45-
// expected-note@+2 {{previous attribute is here}}
46-
cgh.single_task<class T4>([=]() [[intel::max_work_group_size(1, 1, 256),
47-
intel::min_work_groups_per_cu(4),
48-
intel::min_work_groups_per_cu(8)]] {
49-
volatile int A = 42;
50-
});
51-
52-
// expected-error@+1 {{'min_work_groups_per_cu' attribute requires a non-negative integral compile time constant expression}}
53-
cgh.single_task<class T5>([=]() [[intel::min_work_groups_per_cu(-8)]] {
54-
volatile int A = 42;
55-
});
56-
});
22+
// expected-warning@+4 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
23+
cgh.single_task<class T1>(
24+
[=] [[intel::max_work_group_size(1, 1, 256),
25+
intel::min_work_groups_per_cu(2),
26+
intel::max_work_groups_per_mp(4)]] () { volatile int A = 42; });
27+
28+
constexpr float A = 2.0;
29+
// expected-error@+4 {{'min_work_groups_per_cu' attribute requires an integer constant}}
30+
// expected-warning@+4 {{'maxclusterrank' requires sm_90 or higher, CUDA arch provided: sm_70, ignoring 'max_work_groups_per_mp' attribute}}
31+
cgh.single_task<class T2>(
32+
[=] [[intel::max_work_group_size(1, 1, 256),
33+
intel::min_work_groups_per_cu(A),
34+
intel::max_work_groups_per_mp(4)]] () { volatile int A = 42; });
35+
36+
// expected-error@+3 {{'min_work_groups_per_cu' attribute requires an integer constant}}
37+
cgh.single_task<class T3>(
38+
[=] [[intel::max_work_group_size(1, 1, 256),
39+
intel::min_work_groups_per_cu(2147483647 + 1)]] () {
40+
volatile int A = 42;
41+
});
42+
43+
// expected-warning@+5 {{attribute 'min_work_groups_per_cu' is already applied with different arguments}}
44+
// expected-note@+3 {{previous attribute is here}}
45+
cgh.single_task<class T4>(
46+
[=] [[intel::max_work_group_size(1, 1, 256),
47+
intel::min_work_groups_per_cu(4),
48+
intel::min_work_groups_per_cu(8)]] () { volatile int A = 42; });
49+
50+
// expected-error@+2 {{'min_work_groups_per_cu' attribute requires a non-negative integral compile time constant expression}}
51+
cgh.single_task<class T5>(
52+
[=] [[intel::min_work_groups_per_cu(-8)]] () { volatile int A = 42; });
53+
});
5754

5855
Q.submit([&](sycl::handler &cgh) {
5956
cgh.single_task<class F>(Functor<512, 8, 16>{});

‎clang/test/SemaSYCL/lb_sm_90_ast.cpp

Lines changed: 10 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump -triple nvptx-unknown-unknown -target-cpu sm_90 %s | FileCheck %s
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump -triple nvptx-unknown-unknown -target-cpu sm_90 -Wno-c++23-extensions %s | FileCheck %s
22

33
// Tests for AST of Intel max_work_group_size, min_work_groups_per_cu and
44
// max_work_groups_per_mp attribute.
@@ -77,22 +77,20 @@ func1() {}
7777
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6
7878
template <int N>
7979
[[intel::max_work_group_size(N, 8, 8), intel::min_work_groups_per_cu(N),
80-
intel::max_work_groups_per_mp(N)]] void func2() {}
80+
intel::max_work_groups_per_mp(N)]] void
81+
func2() {}
8182

8283
class KernelFunctor {
8384
public:
84-
void operator()() const {
85-
func1();
86-
}
85+
void operator()() const { func1(); }
8786
};
8887

8988
// Test that checks template parameter support on class member function.
90-
template <int N>
91-
class KernelFunctor2 {
89+
template <int N> class KernelFunctor2 {
9290
public:
9391
[[intel::max_work_group_size(N, 8, 8), intel::min_work_groups_per_cu(N),
94-
intel::max_work_groups_per_mp(N)]] void operator()() const {
95-
}
92+
intel::max_work_groups_per_mp(N)]] void
93+
operator()() const {}
9694
};
9795

9896
int main() {
@@ -167,8 +165,9 @@ int main() {
167165
// CHECK-NEXT: value: Int 6
168166
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 6
169167
h.single_task<class kernel_name_3>(
170-
[]() [[intel::max_work_group_size(8, 8, 8), intel::min_work_groups_per_cu(4),
171-
intel::max_work_groups_per_mp(6)]]{});
168+
[] [[intel::max_work_group_size(8, 8, 8),
169+
intel::min_work_groups_per_cu(4),
170+
intel::max_work_groups_per_mp(6)]] () {});
172171
});
173172

174173
func2<6>();

0 commit comments

Comments
 (0)