Skip to content

Commit a17251d

Browse files
schittirschittir
committed
Handle case where attribute parameter is value dependent
Expand test
1 parent ace7f7b commit a17251d

File tree

6 files changed

+91
-29
lines changed

6 files changed

+91
-29
lines changed

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1227,7 +1227,7 @@ def SYCLIntelNumSimdWorkItems : InheritableAttr {
12271227

12281228
def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr {
12291229
let Spellings = [CXX11<"intelfpga","scheduler_target_fmax_mhz">];
1230-
let Args = [UnsignedArgument<"Number">];
1230+
let Args = [ExprArgument<"TargetFmaxMhz">];
12311231
let LangOpts = [SYCLIsDevice, SYCLIsHost];
12321232
let Subjects = SubjectList<[Function], ErrorDiag>;
12331233
let Documentation = [SYCLIntelSchedulerTargetFmaxMhzAttrDocs];

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

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10045,6 +10045,16 @@ class Sema final {
1004510045
bool checkAllowedSYCLInitializer(VarDecl *VD,
1004610046
bool CheckValueDependent = false);
1004710047

10048+
// Adds an intel_reqd_sub_group_size attribute to a particular declaration.
10049+
void addIntelReqdSubGroupSizeAttr(Decl *D, const AttributeCommonInfo &CI,
10050+
Expr *E);
10051+
10052+
// Adds an scheduler_target_fmax_mhz intel_reqd_sub_group_size attribute to a
10053+
// particular declaration.
10054+
void addSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
10055+
const AttributeCommonInfo &CI,
10056+
Expr *E);
10057+
1004810058
//===--------------------------------------------------------------------===//
1004910059
// C++ Coroutines TS
1005010060
//

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -641,8 +641,11 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
641641

642642
if (const SYCLIntelSchedulerTargetFmaxMhzAttr *A =
643643
FD->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
644+
Optional<llvm::APSInt> ArgVal =
645+
A->getTargetFmaxMhz()->getIntegerConstantExpr(FD->getASTContext());
646+
assert(ArgVal.hasValue() && "Not an integer constant expression");
644647
llvm::Metadata *AttrMDArgs[] = {
645-
llvm::ConstantAsMetadata::get(Builder.getInt32(A->getNumber()))};
648+
llvm::ConstantAsMetadata::get(Builder.getInt32(ArgVal->getZExtValue()))};
646649
Fn->setMetadata("scheduler_target_fmax_mhz",
647650
llvm::MDNode::get(Context, AttrMDArgs));
648651
}

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 39 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -3005,32 +3005,48 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
30053005
S.addIntelSYCLSingleArgFunctionAttr<SYCLIntelNumSimdWorkItemsAttr>(D, Attr,
30063006
E);
30073007
}
3008-
// Handles scheduler_target_fmax_mhz
3009-
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3010-
const ParsedAttr &Attr) {
3011-
if (D->isInvalidDecl())
3012-
return;
30133008

3014-
uint32_t TargetFmaxMhz = 0;
3015-
const Expr *E = Attr.getArgAsExpr(0);
3016-
if (!checkUInt32Argument(S, Attr, E, TargetFmaxMhz, 0,
3017-
/*StrictlyUnsigned=*/true))
3009+
// Add scheduler_target_fmax_mhz
3010+
void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr(
3011+
Decl *D, const AttributeCommonInfo &Attr, Expr *E) {
3012+
if (!E)
30183013
return;
30193014

3020-
uint32_t UpperLimit = 1048576;
3021-
if (TargetFmaxMhz > UpperLimit) {
3022-
// Allow frequency to be a "large enough" positive value, zero included
3023-
// Let FPGA backend handle unachievable frequency value
3024-
S.Diag(Attr.getLoc(), diag::err_attribute_argument_out_of_range)
3025-
<< Attr << E->getSourceRange() << 0 << UpperLimit;
3026-
return;
3015+
if (!E->isInstantiationDependent()) {
3016+
Optional<llvm::APSInt> ArgVal = E->getIntegerConstantExpr(getASTContext());
3017+
if (!ArgVal) {
3018+
Diag(E->getExprLoc(), diag::err_attribute_argument_type)
3019+
<< Attr.getAttrName() << AANT_ArgumentIntegerConstant
3020+
<< E->getSourceRange();
3021+
return;
3022+
}
3023+
uint32_t TargetFmaxMhz = ArgVal->getSExtValue();
3024+
uint32_t UpperLimit = 1048576;
3025+
if (TargetFmaxMhz > UpperLimit) {
3026+
// Allow frequency to be a "large enough" positive value, zero included
3027+
// Let FPGA backend handle unachievable frequency value
3028+
Diag(E->getExprLoc(), diag::err_attribute_argument_out_of_range)
3029+
<< Attr.getAttrName() << 0 << UpperLimit << E->getSourceRange();
3030+
return;
3031+
}
30273032
}
30283033

3034+
D->addAttr(::new (Context)
3035+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, Attr, E));
3036+
}
3037+
3038+
// Handle scheduler_target_fmax_mhz
3039+
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3040+
const ParsedAttr &AL) {
3041+
if (D->isInvalidDecl())
3042+
return;
3043+
3044+
Expr *E = AL.getArgAsExpr(0);
3045+
30293046
if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
3030-
S.Diag(Attr.getLoc(), diag::warn_duplicate_attribute) << Attr;
3047+
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
30313048

3032-
D->addAttr(::new (S.Context) SYCLIntelSchedulerTargetFmaxMhzAttr(
3033-
S.Context, Attr, TargetFmaxMhz));
3049+
S.addSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E);
30343050
}
30353051

30363052
// Handles max_global_work_dim.
@@ -8344,6 +8360,10 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
83448360
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
83458361
D->setInvalidDecl();
83468362
}
8363+
} else if (const auto *A =
8364+
D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
8365+
Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
8366+
D->setInvalidDecl();
83478367
} else if (!D->hasAttr<CUDAGlobalAttr>()) {
83488368
if (const auto *A = D->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
83498369
Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -582,6 +582,17 @@ static void instantiateIntelSYCLFunctionAttr(
582582
Result.getAs<Expr>());
583583
}
584584

585+
static void instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
586+
Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
587+
const SYCLIntelSchedulerTargetFmaxMhzAttr *Attr, Decl *New) {
588+
// The TargetFmaxMhz expression is a constant expression.
589+
EnterExpressionEvaluationContext Unevaluated(
590+
S, Sema::ExpressionEvaluationContext::ConstantEvaluated);
591+
ExprResult Result = S.SubstExpr(Attr->getTargetFmaxMhz(), TemplateArgs);
592+
if (!Result.isInvalid())
593+
S.addSYCLIntelSchedulerTargetFmaxMhzAttr(New, *Attr, Result.getAs<Expr>());
594+
}
595+
585596
void Sema::InstantiateAttrsForDecl(
586597
const MultiLevelTemplateArgumentList &TemplateArgs, const Decl *Tmpl,
587598
Decl *New, LateInstantiatedAttrVec *LateAttrs,
@@ -737,6 +748,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
737748
*this, TemplateArgs, SYCLIntelNumSimdWorkItems, New);
738749
continue;
739750
}
751+
if (const auto *SYCLIntelSchedulerTargetFmaxMhz =
752+
dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(TmplAttr)) {
753+
instantiateSYCLIntelSchedulerTargetFmaxMhzAttr(
754+
*this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New);
755+
continue;
756+
}
740757
// Existing DLL attribute on the instantiation takes precedence.
741758
if (TmplAttr->getKind() == attr::DLLExport ||
742759
TmplAttr->getKind() == attr::DLLImport) {

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

Lines changed: 20 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -7,30 +7,42 @@
77
[[intelfpga::scheduler_target_fmax_mhz(2)]] // expected-no-diagnostics
88
void
99
func() {}
10+
11+
template <int N>
12+
[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {}
1013
#endif // TRIGGER_ERROR
1114

1215
int main() {
1316
#ifndef TRIGGER_ERROR
14-
// CHECK-LABEL: -FunctionDecl {{.*}}test_kernel1 'void ()'
15-
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} 5
17+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()'
18+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
19+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
1620
cl::sycl::kernel_single_task<class test_kernel1>(
1721
[]() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});
1822

19-
// CHECK-LABEL: `-FunctionDecl {{.*}}test_kernel2 'void ()'
20-
// CHECK: -SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} 2
23+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()'
24+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
25+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2
2126
cl::sycl::kernel_single_task<class test_kernel2>(
2227
[]() { func(); });
2328

29+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()'
30+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
31+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
32+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
33+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
34+
cl::sycl::kernel_single_task<class test_kernel3>(
35+
[]() { zoo<75>(); });
2436
#else
2537
[[intelfpga::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}
2638

27-
cl::sycl::kernel_single_task<class test_kernel3>(
28-
[]() [[intelfpga::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
29-
3039
cl::sycl::kernel_single_task<class test_kernel4>(
31-
[]() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
40+
[]() [[intelfpga::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
3241

3342
cl::sycl::kernel_single_task<class test_kernel5>(
43+
[]() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
44+
45+
cl::sycl::kernel_single_task<class test_kernel6>(
3446
[]() [[intelfpga::scheduler_target_fmax_mhz(1), intelfpga::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different parameters}}
3547
#endif // TRIGGER_ERROR
3648
}

0 commit comments

Comments
 (0)