Skip to content

Commit 6c15476

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 4f39519 + baed6a5 commit 6c15476

27 files changed

+666
-21
lines changed

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

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1225,6 +1225,24 @@ def SYCLIntelNumSimdWorkItems : InheritableAttr {
12251225
let PragmaAttributeSupport = 0;
12261226
}
12271227

1228+
def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr {
1229+
let Spellings = [CXX11<"intelfpga","scheduler_target_fmax_mhz">];
1230+
let Args = [ExprArgument<"Value">];
1231+
let LangOpts = [SYCLIsDevice, SYCLIsHost];
1232+
let Subjects = SubjectList<[Function], ErrorDiag>;
1233+
let Documentation = [SYCLIntelSchedulerTargetFmaxMhzAttrDocs];
1234+
let PragmaAttributeSupport = 0;
1235+
let AdditionalMembers = [{
1236+
static unsigned getMinValue() {
1237+
return 0;
1238+
}
1239+
static unsigned getMaxValue() {
1240+
return 1048576;
1241+
}
1242+
}];
1243+
1244+
}
1245+
12281246
def SYCLIntelMaxWorkGroupSize : InheritableAttr {
12291247
let Spellings = [CXX11<"intelfpga","max_work_group_size">];
12301248
let Args = [UnsignedArgument<"XDim">,

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2223,6 +2223,28 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
22232223
}];
22242224
}
22252225

2226+
def SYCLIntelSchedulerTargetFmaxMhzAttrDocs : Documentation {
2227+
let Category = DocCatFunction;
2228+
let Heading = "scheduler_target_fmax_mhz (IntelFPGA)";
2229+
let Content = [{
2230+
Applies to a device function/lambda function. Indicates that the kernel should
2231+
be pipelined so as to achieve the specified target clock frequency (Fmax) of N
2232+
MHz. The argument N may be a template parameter. This attribute should be
2233+
ignored for the FPGA emulator device.
2234+
2235+
``[[intelfpga::scheduler_target_fmax_mhz(N)]]``
2236+
Valid values of N are integers in the range [0, 1048576]. The upper limit,
2237+
although too high to be a realistic value for frequency, is chosen to be future
2238+
proof. The FPGA backend emits a diagnostic message if the passed value is
2239+
unachievable by the device.
2240+
2241+
This attribute enables communication of the desired maximum frequency of the
2242+
device operation, guiding the FPGA backend to insert the appropriate number of
2243+
registers to break-up the combinational logic circuit, and thereby controlling
2244+
the length of the longest combinational path.
2245+
}];
2246+
}
2247+
22262248
def SYCLIntelNoGlobalWorkOffsetAttrDocs : Documentation {
22272249
let Category = DocCatFunction;
22282250
let Heading = "no_global_work_offset (IntelFPGA)";

‎clang/include/clang/Basic/AttributeCommonInfo.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -162,6 +162,7 @@ class AttributeCommonInfo {
162162
(ParsedAttr == AT_ReqdWorkGroupSize && isCXX11Attribute()) ||
163163
(ParsedAttr == AT_IntelReqdSubGroupSize && isCXX11Attribute()) ||
164164
ParsedAttr == AT_SYCLIntelNumSimdWorkItems ||
165+
ParsedAttr == AT_SYCLIntelSchedulerTargetFmaxMhz ||
165166
ParsedAttr == AT_SYCLIntelMaxWorkGroupSize ||
166167
ParsedAttr == AT_SYCLIntelMaxGlobalWorkDim ||
167168
ParsedAttr == AT_SYCLIntelNoGlobalWorkOffset)

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

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10062,6 +10062,11 @@ class Sema final {
1006210062
bool checkAllowedSYCLInitializer(VarDecl *VD,
1006310063
bool CheckValueDependent = false);
1006410064

10065+
// Adds a scheduler_target_fmax_mhz attribute to a particular declaration.
10066+
void addSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
10067+
const AttributeCommonInfo &CI,
10068+
Expr *E);
10069+
1006510070
//===--------------------------------------------------------------------===//
1006610071
// C++ Coroutines TS
1006710072
//

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -642,6 +642,17 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
642642
llvm::MDNode::get(Context, AttrMDArgs));
643643
}
644644

645+
if (const SYCLIntelSchedulerTargetFmaxMhzAttr *A =
646+
FD->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
647+
Optional<llvm::APSInt> ArgVal =
648+
A->getValue()->getIntegerConstantExpr(FD->getASTContext());
649+
assert(ArgVal.hasValue() && "Not an integer constant expression");
650+
llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get(
651+
Builder.getInt32(ArgVal->getSExtValue()))};
652+
Fn->setMetadata("scheduler_target_fmax_mhz",
653+
llvm::MDNode::get(Context, AttrMDArgs));
654+
}
655+
645656
if (const SYCLIntelMaxWorkGroupSizeAttr *A =
646657
FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
647658
llvm::Metadata *AttrMDArgs[] = {

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3011,6 +3011,38 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
30113011
E);
30123012
}
30133013

3014+
// Add scheduler_target_fmax_mhz
3015+
void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr(
3016+
Decl *D, const AttributeCommonInfo &Attr, Expr *E) {
3017+
assert(E && "Attribute must have an argument.");
3018+
3019+
SYCLIntelSchedulerTargetFmaxMhzAttr TmpAttr(Context, Attr, E);
3020+
if (!E->isValueDependent()) {
3021+
ExprResult ResultExpr;
3022+
if (checkRangedIntegralArgument<SYCLIntelSchedulerTargetFmaxMhzAttr>(
3023+
E, &TmpAttr, ResultExpr))
3024+
return;
3025+
E = ResultExpr.get();
3026+
}
3027+
3028+
D->addAttr(::new (Context)
3029+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, Attr, E));
3030+
}
3031+
3032+
// Handle scheduler_target_fmax_mhz
3033+
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3034+
const ParsedAttr &AL) {
3035+
if (D->isInvalidDecl())
3036+
return;
3037+
3038+
Expr *E = AL.getArgAsExpr(0);
3039+
3040+
if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
3041+
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
3042+
3043+
S.addSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E);
3044+
}
3045+
30143046
// Handles max_global_work_dim.
30153047
static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D,
30163048
const ParsedAttr &Attr) {
@@ -8230,6 +8262,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
82308262
case ParsedAttr::AT_SYCLIntelNumSimdWorkItems:
82318263
handleNumSimdWorkItemsAttr(S, D, AL);
82328264
break;
8265+
case ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz:
8266+
handleSchedulerTargetFmaxMhzAttr(S, D, AL);
8267+
break;
82338268
case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
82348269
handleMaxGlobalWorkDimAttr(S, D, AL);
82358270
break;

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -531,6 +531,9 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
531531
if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
532532
Attrs.insert(A);
533533

534+
if (auto *A = FD->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
535+
Attrs.insert(A);
536+
534537
if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
535538
Attrs.insert(A);
536539

@@ -3166,6 +3169,7 @@ void Sema::MarkDevice(void) {
31663169
}
31673170
case attr::Kind::SYCLIntelKernelArgsRestrict:
31683171
case attr::Kind::SYCLIntelNumSimdWorkItems:
3172+
case attr::Kind::SYCLIntelSchedulerTargetFmaxMhz:
31693173
case attr::Kind::SYCLIntelMaxGlobalWorkDim:
31703174
case attr::Kind::SYCLIntelNoGlobalWorkOffset:
31713175
case attr::Kind::SYCLSimd: {

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -737,6 +737,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
737737
*this, TemplateArgs, SYCLIntelNumSimdWorkItems, New);
738738
continue;
739739
}
740+
if (const auto *SYCLIntelSchedulerTargetFmaxMhz =
741+
dyn_cast<SYCLIntelSchedulerTargetFmaxMhzAttr>(TmplAttr)) {
742+
instantiateIntelSYCLFunctionAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(
743+
*this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New);
744+
continue;
745+
}
740746
// Existing DLL attribute on the instantiation takes precedence.
741747
if (TmplAttr->getKind() == attr::DLLExport ||
742748
TmplAttr->getKind() == attr::DLLImport) {

‎clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
2+
3+
#include "Inputs/sycl.hpp"
4+
[[intelfpga::scheduler_target_fmax_mhz(5)]] void
5+
func() {}
6+
7+
template <int N>
8+
[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {}
9+
10+
int main() {
11+
cl::sycl::kernel_single_task<class test_kernel1>(
12+
[]() [[intelfpga::scheduler_target_fmax_mhz(2)]]{});
13+
14+
cl::sycl::kernel_single_task<class test_kernel2>(
15+
[]() { func(); });
16+
17+
cl::sycl::kernel_single_task<class test_kernel3>(
18+
[]() { zoo<75>(); });
19+
}
20+
// CHECK: define spir_kernel void @{{.*}}test_kernel1() {{.*}} !scheduler_target_fmax_mhz ![[PARAM1:[0-9]+]]
21+
// CHECK: define spir_kernel void @{{.*}}test_kernel2() {{.*}} !scheduler_target_fmax_mhz ![[PARAM2:[0-9]+]]
22+
// CHECK: define spir_kernel void @{{.*}}test_kernel3() {{.*}} !scheduler_target_fmax_mhz ![[PARAM3:[0-9]+]]
23+
// CHECK: ![[PARAM1]] = !{i32 2}
24+
// CHECK: ![[PARAM2]] = !{i32 5}
25+
// CHECK: ![[PARAM3]] = !{i32 75}

‎clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s
2+
3+
#include "Inputs/sycl.hpp"
4+
[[intelfpga::scheduler_target_fmax_mhz(2)]] void
5+
func() {}
6+
7+
template <int N>
8+
[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {}
9+
10+
int main() {
11+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()'
12+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
13+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
14+
// CHECK-NEXT: value: Int 5
15+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
16+
cl::sycl::kernel_single_task<class test_kernel1>(
17+
[]() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});
18+
19+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()'
20+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
21+
// CHECK-NEXT: ConstantExpr {{.*}} 'int'
22+
// CHECK-NEXT: value: Int 2
23+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2
24+
cl::sycl::kernel_single_task<class test_kernel2>(
25+
[]() { func(); });
26+
27+
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()'
28+
// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}}
29+
// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int'
30+
// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N
31+
// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75
32+
cl::sycl::kernel_single_task<class test_kernel3>(
33+
[]() { zoo<75>(); });
34+
35+
[[intelfpga::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}
36+
37+
cl::sycl::kernel_single_task<class test_kernel4>(
38+
[]() [[intelfpga::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
39+
40+
cl::sycl::kernel_single_task<class test_kernel5>(
41+
[]() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}}
42+
43+
cl::sycl::kernel_single_task<class test_kernel6>(
44+
[]() [[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}}
45+
}

0 commit comments

Comments
 (0)