Skip to content

Commit 5c91609

Browse files
schittirschittir
authored
[SYCL] Propagate attributes from transitive calls to kernel (#1878)
- Remove the (incorrect) distinction between direct and transitive calls for the propagation of attributes to kernel - Handle conflicting attributes - Add a test showing propagation of attributes from transitive call to kernel - Correct failing lit tests to reflect new behavior
1 parent 40100a9 commit 5c91609

File tree

6 files changed

+115
-61
lines changed

6 files changed

+115
-61
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 49 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -510,52 +510,25 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
510510

511511
if (auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
512512
Attrs.insert(A);
513+
513514
if (auto *A = FD->getAttr<ReqdWorkGroupSizeAttr>())
514515
Attrs.insert(A);
515-
// Allow the following kernel attributes only on lambda functions and
516-
// function objects that are called directly from a kernel (i.e. the one
517-
// passed to the parallel_for function). For all other cases,
518-
// emit a warning and ignore.
519-
if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>()) {
520-
if (ParentFD == SYCLKernel) {
521-
Attrs.insert(A);
522-
} else {
523-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
524-
FD->dropAttr<SYCLIntelKernelArgsRestrictAttr>();
525-
}
526-
}
527-
if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>()) {
528-
if (ParentFD == SYCLKernel) {
529-
Attrs.insert(A);
530-
} else {
531-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
532-
FD->dropAttr<SYCLIntelNumSimdWorkItemsAttr>();
533-
}
534-
}
535-
if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
536-
if (ParentFD == SYCLKernel) {
537-
Attrs.insert(A);
538-
} else {
539-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
540-
FD->dropAttr<SYCLIntelMaxWorkGroupSizeAttr>();
541-
}
542-
}
543-
if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
544-
if (ParentFD == SYCLKernel) {
545-
Attrs.insert(A);
546-
} else {
547-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
548-
FD->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
549-
}
550-
}
551-
if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
552-
if (ParentFD == SYCLKernel) {
553-
Attrs.insert(A);
554-
} else {
555-
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
556-
FD->dropAttr<SYCLIntelNoGlobalWorkOffsetAttr>();
557-
}
558-
}
516+
517+
if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>())
518+
Attrs.insert(A);
519+
520+
if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
521+
Attrs.insert(A);
522+
523+
if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
524+
Attrs.insert(A);
525+
526+
if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>())
527+
Attrs.insert(A);
528+
529+
if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>())
530+
Attrs.insert(A);
531+
559532
if (auto *A = FD->getAttr<SYCLSimdAttr>())
560533
Attrs.insert(A);
561534
// Propagate the explicit SIMD attribute through call graph - it is used
@@ -2051,6 +2024,38 @@ void Sema::MarkDevice(void) {
20512024
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
20522025
SYCLKernel->setInvalidDecl();
20532026
}
2027+
} else if (auto *Existing =
2028+
SYCLKernel->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
2029+
if (Existing->getXDim() < Attr->getXDim() ||
2030+
Existing->getYDim() < Attr->getYDim() ||
2031+
Existing->getZDim() < Attr->getZDim()) {
2032+
Diag(SYCLKernel->getLocation(),
2033+
diag::err_conflicting_sycl_kernel_attributes);
2034+
Diag(Existing->getLocation(), diag::note_conflicting_attribute);
2035+
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
2036+
SYCLKernel->setInvalidDecl();
2037+
} else {
2038+
SYCLKernel->addAttr(A);
2039+
}
2040+
} else {
2041+
SYCLKernel->addAttr(A);
2042+
}
2043+
break;
2044+
}
2045+
case attr::Kind::SYCLIntelMaxWorkGroupSize: {
2046+
auto *Attr = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
2047+
if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {
2048+
if (Existing->getXDim() > Attr->getXDim() ||
2049+
Existing->getYDim() > Attr->getYDim() ||
2050+
Existing->getZDim() > Attr->getZDim()) {
2051+
Diag(SYCLKernel->getLocation(),
2052+
diag::err_conflicting_sycl_kernel_attributes);
2053+
Diag(Existing->getLocation(), diag::note_conflicting_attribute);
2054+
Diag(Attr->getLocation(), diag::note_conflicting_attribute);
2055+
SYCLKernel->setInvalidDecl();
2056+
} else {
2057+
SYCLKernel->addAttr(A);
2058+
}
20542059
} else {
20552060
SYCLKernel->addAttr(A);
20562061
}
@@ -2059,7 +2064,6 @@ void Sema::MarkDevice(void) {
20592064
case attr::Kind::SYCLIntelKernelArgsRestrict:
20602065
case attr::Kind::SYCLIntelNumSimdWorkItems:
20612066
case attr::Kind::SYCLIntelMaxGlobalWorkDim:
2062-
case attr::Kind::SYCLIntelMaxWorkGroupSize:
20632067
case attr::Kind::SYCLIntelNoGlobalWorkOffset:
20642068
case attr::Kind::SYCLSimd: {
20652069
if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&

‎clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -verify
2+
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
3+
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
4+
5+
#ifndef TRIGGER_ERROR
6+
[[intelfpga::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics
7+
8+
[[intel::reqd_sub_group_size(1)]] void func_one() {
9+
not_direct_one();
10+
}
11+
12+
#else
13+
[[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note {{conflicting attribute is here}}
14+
15+
[[intelfpga::max_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
16+
void
17+
func_two() {
18+
not_direct_two();
19+
}
20+
21+
[[cl::reqd_work_group_size(4, 4, 4)]] // expected-note 2 {{conflicting attribute is here}}
22+
void
23+
func_three() {
24+
not_direct_two();
25+
}
26+
#endif
27+
28+
template <typename Name, typename Type>
29+
[[clang::sycl_kernel]] void __my_kernel__(Type bar) {
30+
bar();
31+
#ifndef TRIGGER_ERROR
32+
func_one();
33+
#else
34+
func_two();
35+
func_three();
36+
#endif
37+
}
38+
39+
template <typename Name, typename Type>
40+
void parallel_for(Type lambda) {
41+
__my_kernel__<Name>(lambda);
42+
}
43+
44+
void invoke_foo2() {
45+
#ifndef TRIGGER_ERROR
46+
// CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()'
47+
// CHECK: `-FunctionDecl {{.*}}KernelName 'void ()'
48+
// CHECK: -IntelReqdSubGroupSizeAttr {{.*}}
49+
// CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
50+
parallel_for<class KernelName>([]() {});
51+
#else
52+
parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
53+
#endif
54+
}

‎clang/test/SemaSYCL/intel-max-global-work-dim.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::max_global_work_dim(2)]] // expected-warning{{'max_global_work_dim' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::max_global_work_dim(2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::max_global_work_dim(1)]]
@@ -68,9 +67,9 @@ int main() {
6867
[]() [[intelfpga::max_global_work_dim(2)]] {});
6968

7069
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
71-
// CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
70+
// CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}}
7271
kernel<class test_kernel3>(
73-
[]() {func_ignore();});
72+
[]() { func_do_not_ignore(); });
7473

7574
kernel<class test_kernel4>(
7675
TRIFuncObjGood1());

‎clang/test/SemaSYCL/intel-max-work-group-size.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::max_work_group_size(2, 2, 2)]] // expected-warning{{'max_work_group_size' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::max_work_group_size(4, 4, 4)]]
@@ -53,9 +52,9 @@ int main() {
5352
[]() [[intelfpga::max_work_group_size(8, 8, 8)]] {});
5453

5554
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
56-
// CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
55+
// CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}}
5756
kernel<class test_kernel3>(
58-
[]() {func_ignore();});
57+
[]() { func_do_not_ignore(); });
5958

6059
#ifdef TRIGGER_ERROR
6160
[[intelfpga::max_work_group_size(1, 1, 1)]] int Var = 0; // expected-error{{'max_work_group_size' attribute only applies to functions}}

‎clang/test/SemaSYCL/intel-restrict.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,7 @@
11
// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DCHECKDIAG -verify
22
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
33

4-
[[intel::kernel_args_restrict]] // expected-warning{{'kernel_args_restrict' attribute ignored}}
5-
void func_ignore() {}
4+
[[intel::kernel_args_restrict]] void func_do_not_ignore() {}
65

76
struct FuncObj {
87
[[intel::kernel_args_restrict]]
@@ -29,7 +28,7 @@ int main() {
2928
[]() [[intel::kernel_args_restrict]] {});
3029

3130
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
32-
// CHECK-NOT: SYCLIntelKernelArgsRestrictAttr
31+
// CHECK: SYCLIntelKernelArgsRestrictAttr
3332
kernel<class test_kernel3>(
34-
[]() {func_ignore();});
33+
[]() { func_do_not_ignore(); });
3534
}

‎clang/test/SemaSYCL/num_simd_work_items.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -20,8 +20,7 @@ void foo() {
2020

2121
#else // __SYCL_DEVICE_ONLY__
2222

23-
[[intelfpga::num_simd_work_items(2)]] // expected-warning{{'num_simd_work_items' attribute ignored}}
24-
void func_ignore() {}
23+
[[intelfpga::num_simd_work_items(2)]] void func_do_not_ignore() {}
2524

2625
struct FuncObj {
2726
[[intelfpga::num_simd_work_items(42)]]
@@ -45,9 +44,9 @@ int main() {
4544
[]() [[intelfpga::num_simd_work_items(8)]] {});
4645

4746
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
48-
// CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
47+
// CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
4948
kernel<class test_kernel3>(
50-
[]() {func_ignore();});
49+
[]() { func_do_not_ignore(); });
5150

5251
#ifdef TRIGGER_ERROR
5352
[[intelfpga::num_simd_work_items(0)]] int Var = 0; // expected-error{{'num_simd_work_items' attribute only applies to functions}}

0 commit comments

Comments
 (0)