Skip to content

Commit 43ca44a

Browse files
smanna12smanna12
authored
[SYCL] Refactor two FPGA function attributes (#3274)
This patch 1. refactors two function attributes: [[intel::no_global_work_offset()]] and [[intel::scheduler_target_fmax_mhz()]] using #3224 as a template to better fit for community standards. 2. refactors the way we handle duplicate attributes on a given declaration. 3. handles redeclarations or template instantiations properly. 4. adds test Signed-off-by: Soumi Manna <[email protected]>
1 parent b04251d commit 43ca44a

8 files changed

+268
-24
lines changed

‎clang/include/clang/Sema/Sema.h

Lines changed: 11 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10219,6 +10219,16 @@ class Sema final {
1021910219
SYCLIntelNumSimdWorkItemsAttr *
1022010220
MergeSYCLIntelNumSimdWorkItemsAttr(Decl *D,
1022110221
const SYCLIntelNumSimdWorkItemsAttr &A);
10222+
void AddSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
10223+
const AttributeCommonInfo &CI,
10224+
Expr *E);
10225+
SYCLIntelSchedulerTargetFmaxMhzAttr *MergeSYCLIntelSchedulerTargetFmaxMhzAttr(
10226+
Decl *D, const SYCLIntelSchedulerTargetFmaxMhzAttr &A);
10227+
void AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D,
10228+
const AttributeCommonInfo &CI,
10229+
Expr *E);
10230+
SYCLIntelNoGlobalWorkOffsetAttr *MergeSYCLIntelNoGlobalWorkOffsetAttr(
10231+
Decl *D, const SYCLIntelNoGlobalWorkOffsetAttr &A);
1022210232

1022310233
/// AddAlignedAttr - Adds an aligned attribute to a particular declaration.
1022410234
void AddAlignedAttr(Decl *D, const AttributeCommonInfo &CI, Expr *E,
@@ -13099,8 +13109,7 @@ void Sema::addIntelSingleArgAttr(Decl *D, const AttributeCommonInfo &CI,
1309913109
return;
1310013110
}
1310113111
}
13102-
if (CI.getParsedKind() == ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz ||
13103-
CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) {
13112+
if (CI.getParsedKind() == ParsedAttr::AT_IntelFPGAPrivateCopies) {
1310413113
if (ArgInt < 0) {
1310513114
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
1310613115
<< CI << /*non-negative*/ 1;

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2622,6 +2622,10 @@ static bool mergeDeclAttribute(Sema &S, NamedDecl *D,
26222622
NewAttr = S.MergeIntelReqdSubGroupSizeAttr(D, *A);
26232623
else if (const auto *A = dyn_cast<SYCLIntelNumSimdWorkItemsAttr>(Attr))
26242624
NewAttr = S.MergeSYCLIntelNumSimdWorkItemsAttr(D, *A);
2625+
else if (const auto *A = dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(Attr))
2626+
NewAttr = S.MergeSYCLIntelSchedulerTargetFmaxMhzAttr(D, *A);
2627+
else if (const auto *A = dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(Attr))
2628+
NewAttr = S.MergeSYCLIntelNoGlobalWorkOffsetAttr(D, *A);
26252629
else if (Attr->shouldInheritEvenIfAlreadyPresent() || !DeclHasAttr(D, Attr))
26262630
NewAttr = cast<InheritableAttr>(Attr->clone(S.Context));
26272631

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 115 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -3366,19 +3366,71 @@ static void handleUseStallEnableClustersAttr(Sema &S, Decl *D,
33663366
}
33673367

33683368
// Handle scheduler_target_fmax_mhz
3369-
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3370-
const ParsedAttr &AL) {
3371-
if (D->isInvalidDecl())
3372-
return;
3369+
void Sema::AddSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
3370+
const AttributeCommonInfo &CI,
3371+
Expr *E) {
3372+
if (!E->isValueDependent()) {
3373+
// Validate that we have an integer constant expression and then store the
3374+
// converted constant expression into the semantic attribute so that we
3375+
// don't have to evaluate it again later.
3376+
llvm::APSInt ArgVal;
3377+
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
3378+
if (Res.isInvalid())
3379+
return;
3380+
E = Res.get();
33733381

3374-
Expr *E = AL.getArgAsExpr(0);
3382+
// This attribute requires a non-negative value.
3383+
if (ArgVal < 0) {
3384+
Diag(E->getExprLoc(), diag::err_attribute_requires_positive_integer)
3385+
<< CI << /*non-negative*/ 1;
3386+
return;
3387+
}
3388+
// Check to see if there's a duplicate attribute with different values
3389+
// already applied to the declaration.
3390+
if (const auto *DeclAttr =
3391+
D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
3392+
// If the other attribute argument is instantiation dependent, we won't
3393+
// have converted it to a constant expression yet and thus we test
3394+
// whether this is a null pointer.
3395+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
3396+
if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
3397+
Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
3398+
Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
3399+
return;
3400+
}
3401+
}
3402+
}
33753403

3376-
if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
3377-
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
3404+
D->addAttr(::new (Context)
3405+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, CI, E));
3406+
}
33783407

3408+
SYCLIntelSchedulerTargetFmaxMhzAttr *
3409+
Sema::MergeSYCLIntelSchedulerTargetFmaxMhzAttr(
3410+
Decl *D, const SYCLIntelSchedulerTargetFmaxMhzAttr &A) {
3411+
// Check to see if there's a duplicate attribute with different values
3412+
// already applied to the declaration.
3413+
if (const auto *DeclAttr =
3414+
D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
3415+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
3416+
const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
3417+
if (DeclExpr && MergeExpr &&
3418+
DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
3419+
Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
3420+
Diag(A.getLoc(), diag::note_previous_attribute);
3421+
return nullptr;
3422+
}
3423+
}
3424+
return ::new (Context)
3425+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, A, A.getValue());
3426+
}
3427+
3428+
static void handleSYCLIntelSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3429+
const ParsedAttr &AL) {
33793430
S.CheckDeprecatedSYCLAttributeSpelling(AL);
33803431

3381-
S.addIntelSingleArgAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(D, AL, E);
3432+
Expr *E = AL.getArgAsExpr(0);
3433+
S.AddSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E);
33823434
}
33833435

33843436
// Handles max_global_work_dim.
@@ -5708,17 +5760,66 @@ static bool checkForDuplicateAttribute(Sema &S, Decl *D,
57085760
return false;
57095761
}
57105762

5711-
static void handleNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
5712-
const ParsedAttr &A) {
5713-
checkForDuplicateAttribute<SYCLIntelNoGlobalWorkOffsetAttr>(S, D, A);
5763+
void Sema::AddSYCLIntelNoGlobalWorkOffsetAttr(Decl *D,
5764+
const AttributeCommonInfo &CI,
5765+
Expr *E) {
5766+
if (!E->isValueDependent()) {
5767+
// Validate that we have an integer constant expression and then store the
5768+
// converted constant expression into the semantic attribute so that we
5769+
// don't have to evaluate it again later.
5770+
llvm::APSInt ArgVal;
5771+
ExprResult Res = VerifyIntegerConstantExpression(E, &ArgVal);
5772+
if (Res.isInvalid())
5773+
return;
5774+
E = Res.get();
5775+
5776+
// Check to see if there's a duplicate attribute with different values
5777+
// already applied to the declaration.
5778+
if (const auto *DeclAttr = D->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
5779+
// If the other attribute argument is instantiation dependent, we won't
5780+
// have converted it to a constant expression yet and thus we test
5781+
// whether this is a null pointer.
5782+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
5783+
if (DeclExpr && ArgVal != DeclExpr->getResultAsAPSInt()) {
5784+
Diag(CI.getLoc(), diag::warn_duplicate_attribute) << CI;
5785+
Diag(DeclAttr->getLoc(), diag::note_previous_attribute);
5786+
return;
5787+
}
5788+
}
5789+
}
5790+
5791+
D->addAttr(::new (Context) SYCLIntelNoGlobalWorkOffsetAttr(Context, CI, E));
5792+
}
5793+
5794+
SYCLIntelNoGlobalWorkOffsetAttr *Sema::MergeSYCLIntelNoGlobalWorkOffsetAttr(
5795+
Decl *D, const SYCLIntelNoGlobalWorkOffsetAttr &A) {
5796+
// Check to see if there's a duplicate attribute with different values
5797+
// already applied to the declaration.
5798+
if (const auto *DeclAttr = D->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
5799+
const auto *DeclExpr = dyn_cast<ConstantExpr>(DeclAttr->getValue());
5800+
const auto *MergeExpr = dyn_cast<ConstantExpr>(A.getValue());
5801+
if (DeclExpr && MergeExpr &&
5802+
DeclExpr->getResultAsAPSInt() != MergeExpr->getResultAsAPSInt()) {
5803+
Diag(DeclAttr->getLoc(), diag::warn_duplicate_attribute) << &A;
5804+
Diag(A.getLoc(), diag::note_previous_attribute);
5805+
return nullptr;
5806+
}
5807+
}
5808+
return ::new (Context)
5809+
SYCLIntelNoGlobalWorkOffsetAttr(Context, A, A.getValue());
5810+
}
5811+
5812+
static void handleSYCLIntelNoGlobalWorkOffsetAttr(Sema &S, Decl *D,
5813+
const ParsedAttr &A) {
57145814
S.CheckDeprecatedSYCLAttributeSpelling(A);
57155815

57165816
// If no attribute argument is specified, set to default value '1'.
57175817
Expr *E = A.isArgExpr(0)
57185818
? A.getArgAsExpr(0)
57195819
: IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
57205820
S.Context.IntTy, A.getLoc());
5721-
S.addIntelSingleArgAttr<SYCLIntelNoGlobalWorkOffsetAttr>(D, A, E);
5821+
5822+
S.AddSYCLIntelNoGlobalWorkOffsetAttr(D, A, E);
57225823
}
57235824

57245825
/// Handle the [[intelfpga::doublepump]] and [[intelfpga::singlepump]] attributes.
@@ -8897,13 +8998,13 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
88978998
handleSYCLIntelNumSimdWorkItemsAttr(S, D, AL);
88988999
break;
88999000
case ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz:
8900-
handleSchedulerTargetFmaxMhzAttr(S, D, AL);
9001+
handleSYCLIntelSchedulerTargetFmaxMhzAttr(S, D, AL);
89019002
break;
89029003
case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
89039004
handleMaxGlobalWorkDimAttr(S, D, AL);
89049005
break;
89059006
case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
8906-
handleNoGlobalWorkOffsetAttr(S, D, AL);
9007+
handleSYCLIntelNoGlobalWorkOffsetAttr(S, D, AL);
89079008
break;
89089009
case ParsedAttr::AT_SYCLIntelUseStallEnableClusters:
89099010
handleUseStallEnableClustersAttr(S, D, AL);

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 22 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -663,6 +663,26 @@ static void instantiateSYCLIntelNumSimdWorkItemsAttr(
663663
S.AddSYCLIntelNumSimdWorkItemsAttr(New, *A, Result.getAs<Expr>());
664664
}
665665

666+
static void instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
667+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
668+
const SYCLIntelSchedulerTargetFmaxMhzAttr *A, Decl *New) {
669+
EnterExpressionEvaluationContext Unevaluated(
670+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
671+
ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs);
672+
if (!Result.isInvalid())
673+
S.AddSYCLIntelSchedulerTargetFmaxMhzAttr(New, *A, Result.getAs<Expr>());
674+
}
675+
676+
static void instantiateSYCLIntelNoGlobalWorkOffsetAttr(
677+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
678+
const SYCLIntelNoGlobalWorkOffsetAttr *A, Decl *New) {
679+
EnterExpressionEvaluationContext Unevaluated(
680+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
681+
ExprResult Result = S.SubstExpr(A->getValue(), TemplateArgs);
682+
if (!Result.isInvalid())
683+
S.AddSYCLIntelNoGlobalWorkOffsetAttr(New, *A, Result.getAs<Expr>());
684+
}
685+
666686
template <typename AttrName>
667687
static void instantiateIntelSYCLFunctionAttr(
668688
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
@@ -866,7 +886,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
866886
}
867887
if (const auto *SYCLIntelSchedulerTargetFmaxMhz =
868888
dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(TmplAttr)) {
869-
instantiateIntelSYCLFunctionAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(
889+
instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
870890
*this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New);
871891
continue;
872892
}
@@ -884,7 +904,7 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
884904
}
885905
if (const auto *SYCLIntelNoGlobalWorkOffset =
886906
dyn_cast<SYCLIntelNoGlobalWorkOffsetAttr>(TmplAttr)) {
887-
instantiateIntelSYCLFunctionAttr<SYCLIntelNoGlobalWorkOffsetAttr>(
907+
instantiateSYCLIntelNoGlobalWorkOffsetAttr(
888908
*this, TemplateArgs, SYCLIntelNoGlobalWorkOffset, New);
889909
continue;
890910
}

‎clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -50,9 +50,9 @@ int main() {
5050
[[intel::no_global_work_offset(1)]] int a;
5151
});
5252

53-
// expected-warning@+2{{attribute 'no_global_work_offset' is already applied}}
5453
h.single_task<class test_kernel7>(
55-
[]() [[intel::no_global_work_offset(0), intel::no_global_work_offset(1)]]{});
54+
[]() [[intel::no_global_work_offset(0), // expected-note {{previous attribute is here}}
55+
intel::no_global_work_offset(1)]]{}); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
5656
});
5757
return 0;
5858
}

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,12 @@
66
[[intelfpga::scheduler_target_fmax_mhz(2)]] void
77
func() {}
88

9+
[[intel::scheduler_target_fmax_mhz(12)]] void bar();
10+
[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK
11+
12+
[[intel::scheduler_target_fmax_mhz(12)]] void baz(); // expected-note {{previous attribute is here}}
13+
[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
14+
915
template <int N>
1016
[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}
1117

@@ -47,5 +53,6 @@ int main() {
4753
[]() [[intel::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
4854

4955
cl::sycl::kernel_single_task<class test_kernel6>(
50-
[]() [[intel::scheduler_target_fmax_mhz(1), intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
56+
[]() [[intel::scheduler_target_fmax_mhz(1), // expected-note {{previous attribute is here}}
57+
intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
5158
}

‎clang/test/SemaSYCL/sycl-device-intel-fpga-no-global-work-offset.cpp

Lines changed: 21 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,18 @@ int main() {
3636
KernelFunctor<1>();
3737
}
3838

39+
[[intel::no_global_work_offset]] void func3 ();
40+
[[intel::no_global_work_offset(1)]] void func3() {} // OK
41+
42+
[[intel::no_global_work_offset(0)]] void func4(); // expected-note {{previous attribute is here}}
43+
[[intel::no_global_work_offset]] void func4(); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
44+
45+
[[intel::no_global_work_offset(1)]] void func5();
46+
[[intel::no_global_work_offset(1)]] void func5() {} // OK
47+
48+
[[intel::no_global_work_offset(0)]] void func6(); // expected-note {{previous attribute is here}}
49+
[[intel::no_global_work_offset(1)]] void func6(); // expected-warning{{attribute 'no_global_work_offset' is already applied with different parameters}}
50+
3951
// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor
4052
// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition
4153
// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor
@@ -48,14 +60,20 @@ int main() {
4860

4961
// Test that checks template parameter suppport on function.
5062
template <int N>
51-
[[intel::no_global_work_offset(N)]] void func3() {}
63+
[[intel::no_global_work_offset(N)]] void func6() {}
64+
65+
template <int N>
66+
[[intel::no_global_work_offset(0)]] void func7(); // expected-note {{previous attribute is here}}
67+
template <int N>
68+
[[intel::no_global_work_offset(N)]] void func7() {} // expected-warning {{attribute 'no_global_work_offset' is already applied with different parameters}}
5269

5370
int check() {
54-
func3<1>();
71+
func6<1>();
72+
func7<1>(); //expected-note {{in instantiation of function template specialization 'func7<1>' requested here}}
5573
return 0;
5674
}
5775

58-
// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()'
76+
// CHECK: FunctionDecl {{.*}} {{.*}} func6 'void ()'
5977
// CHECK: TemplateArgument integral 1
6078
// CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
6179
// CHECK-NEXT: ConstantExpr {{.*}} 'int'

0 commit comments

Comments
 (0)