Skip to content

Commit 20013e2

Browse files
schittirschittir
authored
[SYCL] Add clang support for FPGA kernel attribute scheduler_target_fmax_mhz (intel#2511)
1 parent 73e957f commit 20013e2

File tree

10 files changed

+172
-0
lines changed

10 files changed

+172
-0
lines changed

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

+18
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

+22
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

+1
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

+5
Original file line numberDiff line numberDiff line change
@@ -10060,6 +10060,11 @@ class Sema final {
1006010060
bool checkAllowedSYCLInitializer(VarDecl *VD,
1006110061
bool CheckValueDependent = false);
1006210062

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

‎clang/lib/CodeGen/CodeGenFunction.cpp

+11
Original file line numberDiff line numberDiff line change
@@ -641,6 +641,17 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
641641
llvm::MDNode::get(Context, AttrMDArgs));
642642
}
643643

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

‎clang/lib/Sema/SemaDeclAttr.cpp

+35
Original file line numberDiff line numberDiff line change
@@ -3006,6 +3006,38 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
30063006
E);
30073007
}
30083008

3009+
// Add scheduler_target_fmax_mhz
3010+
void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr(
3011+
Decl *D, const AttributeCommonInfo &Attr, Expr *E) {
3012+
assert(E && "Attribute must have an argument.");
3013+
3014+
SYCLIntelSchedulerTargetFmaxMhzAttr TmpAttr(Context, Attr, E);
3015+
if (!E->isValueDependent()) {
3016+
ExprResult ResultExpr;
3017+
if (checkRangedIntegralArgument<SYCLIntelSchedulerTargetFmaxMhzAttr>(
3018+
E, &TmpAttr, ResultExpr))
3019+
return;
3020+
E = ResultExpr.get();
3021+
}
3022+
3023+
D->addAttr(::new (Context)
3024+
SYCLIntelSchedulerTargetFmaxMhzAttr(Context, Attr, E));
3025+
}
3026+
3027+
// Handle scheduler_target_fmax_mhz
3028+
static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
3029+
const ParsedAttr &AL) {
3030+
if (D->isInvalidDecl())
3031+
return;
3032+
3033+
Expr *E = AL.getArgAsExpr(0);
3034+
3035+
if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
3036+
S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL;
3037+
3038+
S.addSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E);
3039+
}
3040+
30093041
// Handles max_global_work_dim.
30103042
static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D,
30113043
const ParsedAttr &Attr) {
@@ -8225,6 +8257,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
82258257
case ParsedAttr::AT_SYCLIntelNumSimdWorkItems:
82268258
handleNumSimdWorkItemsAttr(S, D, AL);
82278259
break;
8260+
case ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz:
8261+
handleSchedulerTargetFmaxMhzAttr(S, D, AL);
8262+
break;
82288263
case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim:
82298264
handleMaxGlobalWorkDimAttr(S, D, AL);
82308265
break;

‎clang/lib/Sema/SemaSYCL.cpp

+4
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

+6
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

+25
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

+45
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)