Skip to content

Commit 7f37250

Browse files
smanna12smanna12
authored
[FPGA][SYCL] Fix max_work_group_size and reqd_work_group_size attribute arguments check (#5592)
If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]] attribute, this patch checks if values of reqd_work_group_size arguments are equal or less than values of max_work_group_size attribute arguments. Some of the test cases were missed during refactoring work with PGA function attribute [[intel::max_work_group_size()]] on #5392 This patch adds the missing cases and fixes pulldown bug. Signed-off-by: Soumi Manna <[email protected]>
1 parent a9ad3af commit 7f37250

File tree

7 files changed

+182
-26
lines changed

7 files changed

+182
-26
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11597,7 +11597,7 @@ def err_sycl_non_constant_array_type : Error<
1159711597
def err_conflicting_sycl_kernel_attributes : Error<
1159811598
"conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function">;
1159911599
def err_conflicting_sycl_function_attributes : Error<
11600-
"%0 attribute conflicts with '%1' attribute">;
11600+
"%0 attribute conflicts with %1 attribute">;
1160111601
def err_sycl_function_attribute_mismatch : Error<
1160211602
"SYCL kernel without %0 attribute can't call a function with this attribute">;
1160311603
def err_sycl_x_y_z_arguments_must_be_one : Error<

‎clang/lib/Sema/SemaDeclAttr.cpp‎

Lines changed: 110 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -3268,12 +3268,31 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
32683268

32693269
ASTContext &Ctx = S.getASTContext();
32703270

3271+
// The arguments to reqd_work_group_size are ordered based on which index
3272+
// increments the fastest. In OpenCL, the first argument is the index that
3273+
// increments the fastest, and in SYCL, the last argument is the index that
3274+
// increments the fastest.
3275+
//
3276+
// [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
3277+
// available in SYCL modes and follow the SYCL rules.
3278+
// __attribute__((reqd_work_group_size)) is only available in OpenCL mode
3279+
// and follows the OpenCL rules.
32713280
if (const auto *A = D->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
3272-
if (!((getExprValue(AL.getArgAsExpr(0), Ctx) <= *A->getXDimVal()) &&
3273-
(getExprValue(AL.getArgAsExpr(1), Ctx) <= *A->getYDimVal()) &&
3274-
(getExprValue(AL.getArgAsExpr(2), Ctx) <= *A->getZDimVal()))) {
3281+
bool CheckFirstArgument =
3282+
S.getLangOpts().OpenCL
3283+
? getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getZDimVal()
3284+
: getExprValue(AL.getArgAsExpr(0), Ctx) > *A->getXDimVal();
3285+
bool CheckSecondArgument =
3286+
getExprValue(AL.getArgAsExpr(1), Ctx) > *A->getYDimVal();
3287+
bool CheckThirdArgument =
3288+
S.getLangOpts().OpenCL
3289+
? getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getXDimVal()
3290+
: getExprValue(AL.getArgAsExpr(2), Ctx) > *A->getZDimVal();
3291+
3292+
if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
32753293
S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
3276-
<< AL << A->getSpelling();
3294+
<< AL << A;
3295+
S.Diag(A->getLocation(), diag::note_conflicting_attribute);
32773296
Result &= false;
32783297
}
32793298
}
@@ -3286,7 +3305,8 @@ static bool checkWorkGroupSizeValues(Sema &S, Decl *D, const ParsedAttr &AL) {
32863305
(getExprValue(AL.getArgAsExpr(2), Ctx) >=
32873306
getExprValue(A->getZDim(), Ctx)))) {
32883307
S.Diag(AL.getLoc(), diag::err_conflicting_sycl_function_attributes)
3289-
<< AL << A->getSpelling();
3308+
<< AL << A;
3309+
S.Diag(A->getLocation(), diag::note_conflicting_attribute);
32903310
Result &= false;
32913311
}
32923312
}
@@ -3562,6 +3582,23 @@ static bool InvalidWorkGroupSizeAttrs(const Expr *MGValue, const Expr *XDim,
35623582
ZDimExpr->getResultAsAPSInt() != 1));
35633583
}
35643584

3585+
// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
3586+
// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
3587+
// attribute, check to see if values of reqd_work_group_size arguments are
3588+
// equal or less than values of max_work_group_size attribute arguments.
3589+
static bool checkWorkGroupSizeAttrValues(const Expr *RWGS, const Expr *MWGS) {
3590+
// If any of the operand is still value dependent, we can't test anything.
3591+
const auto *RWGSCE = dyn_cast<ConstantExpr>(RWGS);
3592+
const auto *MWGSCE = dyn_cast<ConstantExpr>(MWGS);
3593+
3594+
if (!RWGSCE || !MWGSCE)
3595+
return false;
3596+
3597+
// Otherwise, check if value of reqd_work_group_size argument is
3598+
// greater than value of max_work_group_size attribute argument.
3599+
return RWGSCE->getResultAsAPSInt() > MWGSCE->getResultAsAPSInt();
3600+
}
3601+
35653602
void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
35663603
const AttributeCommonInfo &CI,
35673604
Expr *XDim, Expr *YDim,
@@ -3595,6 +3632,40 @@ void Sema::AddSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
35953632
if (!XDim || !YDim || !ZDim)
35963633
return;
35973634

3635+
// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
3636+
// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
3637+
// attribute, check to see if values of reqd_work_group_size arguments are
3638+
// equal or less than values of max_work_group_size attribute arguments.
3639+
//
3640+
// The arguments to reqd_work_group_size are ordered based on which index
3641+
// increments the fastest. In OpenCL, the first argument is the index that
3642+
// increments the fastest, and in SYCL, the last argument is the index that
3643+
// increments the fastest.
3644+
//
3645+
// [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
3646+
// available in SYCL modes and follow the SYCL rules.
3647+
// __attribute__((reqd_work_group_size)) is only available in OpenCL mode
3648+
// and follows the OpenCL rules.
3649+
if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
3650+
bool CheckFirstArgument =
3651+
getLangOpts().OpenCL
3652+
? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), ZDim)
3653+
: checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), XDim);
3654+
bool CheckSecondArgument =
3655+
checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), YDim);
3656+
bool CheckThirdArgument =
3657+
getLangOpts().OpenCL
3658+
? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), XDim)
3659+
: checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), ZDim);
3660+
3661+
if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
3662+
Diag(CI.getLoc(), diag::err_conflicting_sycl_function_attributes)
3663+
<< CI << DeclAttr;
3664+
Diag(DeclAttr->getLoc(), diag::note_conflicting_attribute);
3665+
return;
3666+
}
3667+
}
3668+
35983669
// If the declaration has a SYCLIntelMaxWorkGroupSizeAttr, check to see if
35993670
// the attribute holds equal values to (1, 1, 1) in case the value of
36003671
// SYCLIntelMaxGlobalWorkDimAttr equals to 0.
@@ -3655,6 +3726,40 @@ SYCLIntelMaxWorkGroupSizeAttr *Sema::MergeSYCLIntelMaxWorkGroupSizeAttr(
36553726
return nullptr;
36563727
}
36573728

3729+
// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
3730+
// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
3731+
// attribute, check to see if values of reqd_work_group_size arguments are
3732+
// equal or less than values of max_work_group_size attribute arguments.
3733+
//
3734+
// The arguments to reqd_work_group_size are ordered based on which index
3735+
// increments the fastest. In OpenCL, the first argument is the index that
3736+
// increments the fastest, and in SYCL, the last argument is the index that
3737+
// increments the fastest.
3738+
//
3739+
// [[sycl::reqd_work_group_size]] and [[cl::reqd_work_group_size]] are
3740+
// available in SYCL modes and follow the SYCL rules.
3741+
// __attribute__((reqd_work_group_size)) is only available in OpenCL mode
3742+
// and follows the OpenCL rules.
3743+
if (const auto *DeclAttr = D->getAttr<ReqdWorkGroupSizeAttr>()) {
3744+
bool CheckFirstArgument =
3745+
getLangOpts().OpenCL
3746+
? checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getZDim())
3747+
: checkWorkGroupSizeAttrValues(DeclAttr->getXDim(), A.getXDim());
3748+
bool CheckSecondArgument =
3749+
checkWorkGroupSizeAttrValues(DeclAttr->getYDim(), A.getYDim());
3750+
bool CheckThirdArgument =
3751+
getLangOpts().OpenCL
3752+
? checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getXDim())
3753+
: checkWorkGroupSizeAttrValues(DeclAttr->getZDim(), A.getZDim());
3754+
3755+
if (CheckFirstArgument || CheckSecondArgument || CheckThirdArgument) {
3756+
Diag(DeclAttr->getLoc(), diag::err_conflicting_sycl_function_attributes)
3757+
<< DeclAttr << &A;
3758+
Diag(A.getLoc(), diag::note_conflicting_attribute);
3759+
return nullptr;
3760+
}
3761+
}
3762+
36583763
// If the declaration has a SYCLIntelMaxWorkGroupSizeAttr,
36593764
// check to see if the attribute holds equal values to
36603765
// (1, 1, 1) in case the value of SYCLIntelMaxGlobalWorkDimAttr

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

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -35,12 +35,14 @@ struct Func {
3535

3636
#ifdef TRIGGER_ERROR
3737
struct DAFuncObj {
38-
[[intel::max_work_group_size(4, 4, 4)]]
39-
[[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
40-
// expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
41-
// expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
42-
void operator()() const {}
38+
[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
39+
[[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \
40+
// expected-warning{{attribute 'cl::reqd_work_group_size' is deprecated}} \
41+
// expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}}
42+
void
43+
operator()() const {}
4344
};
45+
4446
#endif // TRIGGER_ERROR
4547

4648
int main() {

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

Lines changed: 50 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,11 @@ class Functor {
3030
[[intel::max_work_group_size(16, 16, 16)]] [[intel::max_work_group_size(32, 32, 32)]] void operator()(int) const; // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} expected-note {{previous attribute is here}}
3131
};
3232

33+
class FunctorC {
34+
public:
35+
[[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(64, 64, 64)]] void operator()() const;
36+
[[sycl::reqd_work_group_size(64, 64, 64)]] [[intel::max_work_group_size(16, 16, 16)]] void operator()(int) const; // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} expected-note {{conflicting attribute is here}}
37+
};
3338
// Ensure that template arguments behave appropriately based on instantiations.
3439
template <int N>
3540
[[intel::max_work_group_size(N, 1, 1)]] void f6(); // #f6
@@ -59,3 +64,48 @@ void instantiate() {
5964
// expected-note@#f7prev {{previous attribute is here}}
6065
f7<2, 2, 2>(); // expected-note {{in instantiation}}
6166
}
67+
68+
// If the [[intel::max_work_group_size(X, Y, Z)]] attribute is specified on
69+
// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
70+
// attribute, check to see if values of reqd_work_group_size arguments are
71+
// equal or less than values coming from max_work_group_size attribute.
72+
[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
73+
[[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
74+
void
75+
f9() {}
76+
77+
[[intel::max_work_group_size(4, 4, 4)]] void f10();
78+
[[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK
79+
80+
[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK
81+
82+
// FIXME: We do not have support yet for checking
83+
// reqd_work_group_size and max_work_group_size
84+
// attributes when merging, so the test compiles without
85+
// any diagnostic when it shouldn't.
86+
[[sycl::reqd_work_group_size(64, 64, 64)]] void f12();
87+
[[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected error but now OK.
88+
89+
[[intel::max_work_group_size(16, 16, 16)]] // expected-note {{conflicting attribute is here}}
90+
[[sycl::reqd_work_group_size(16, 64, 16)]] void // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
91+
f13() {}
92+
93+
[[intel::max_work_group_size(16, 16, 16)]] void f14(); // expected-note {{conflicting attribute is here}}
94+
[[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
95+
96+
[[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
97+
// expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
98+
[[intel::max_work_group_size(1, 2, 3)]] void
99+
f15() {} // OK
100+
101+
[[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}}
102+
[[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
103+
104+
[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK
105+
106+
[[sycl::reqd_work_group_size(16)]] // expected-note {{conflicting attribute is here}}
107+
[[intel::max_work_group_size(1, 1, 16)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
108+
f18();
109+
110+
[[intel::max_work_group_size(16, 16, 1)]] void f19();
111+
[[sycl::reqd_work_group_size(16, 16)]] void f19(); // OK

‎clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp‎

Lines changed: 4 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -20,13 +20,8 @@ func1();
2020

2121
#else
2222
//second case - expect error
23-
[[intel::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
24-
void
25-
func2();
26-
27-
[[sycl::reqd_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}}
28-
void
29-
func2() {}
23+
[[intel::max_work_group_size(4, 4, 4)]] void func2(); // expected-note {{conflicting attribute is here}}
24+
[[sycl::reqd_work_group_size(8, 8, 8)]] void func2() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}
3025

3126
//third case - expect error
3227
[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
@@ -36,7 +31,7 @@ func3();
3631
[[sycl::reqd_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
3732
void
3833
// expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
39-
func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with ''reqd_work_group_size'' attribute}}
34+
func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
4035

4136
// fourth case - expect warning.
4237
[[intel::max_work_group_size(4, 4, 4)]] void func4(); // expected-note {{previous attribute is here}}
@@ -77,7 +72,7 @@ int main() {
7772

7873
#else
7974
h.single_task<class test_kernel2>(
80-
[]() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
75+
[]() { func2(); });
8176

8277
h.single_task<class test_kernel3>(
8378
[]() { func3(); });

‎clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -26,9 +26,11 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro
2626
class Functor32 {
2727
public:
2828
// expected-note@+3{{conflicting attribute is here}}
29-
// expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}}
30-
// expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
31-
[[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
29+
// expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}}
30+
// expected-error@+2{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
31+
[[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}}
32+
[[sycl::reqd_work_group_size(1, 1, 32)]] void
33+
operator()() const {}
3234
};
3335
#endif // TRIGGER_ERROR
3436

‎clang/test/SemaSYCL/reqd-work-group-size-device.cpp‎

Lines changed: 5 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -49,9 +49,11 @@ class Functor16 {
4949
class Functor32 {
5050
public:
5151
// expected-note@+3{{conflicting attribute is here}}
52-
// expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}}
53-
// expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
54-
[[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {}
52+
// expected-warning@+3{{attribute 'reqd_work_group_size' is already applied with different arguments}}
53+
// expected-error@+2 {{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
54+
[[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{conflicting attribute is here}}
55+
[[sycl::reqd_work_group_size(1, 1, 32)]] void
56+
operator()() const {}
5557
};
5658
#endif
5759
class Functor16x16x16 {

0 commit comments

Comments
 (0)