diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index c96adf49e2c26..0059733872ce4 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1225,6 +1225,24 @@ def SYCLIntelNumSimdWorkItems : InheritableAttr { let PragmaAttributeSupport = 0; } +def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr { + let Spellings = [CXX11<"intelfpga","scheduler_target_fmax_mhz">]; + let Args = [ExprArgument<"Value">]; + let LangOpts = [SYCLIsDevice, SYCLIsHost]; + let Subjects = SubjectList<[Function], ErrorDiag>; + let Documentation = [SYCLIntelSchedulerTargetFmaxMhzAttrDocs]; + let PragmaAttributeSupport = 0; + let AdditionalMembers = [{ + static unsigned getMinValue() { + return 0; + } + static unsigned getMaxValue() { + return 1048576; + } + }]; + +} + def SYCLIntelMaxWorkGroupSize : InheritableAttr { let Spellings = [CXX11<"intelfpga","max_work_group_size">]; let Args = [UnsignedArgument<"XDim">, diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 970b49ccd42fe..e93bfbeae8c94 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2197,6 +2197,28 @@ device kernel, the attribute is ignored and it is not propagated to a kernel. }]; } +def SYCLIntelSchedulerTargetFmaxMhzAttrDocs : Documentation { + let Category = DocCatFunction; + let Heading = "scheduler_target_fmax_mhz (IntelFPGA)"; + let Content = [{ +Applies to a device function/lambda function. Indicates that the kernel should +be pipelined so as to achieve the specified target clock frequency (Fmax) of N +MHz. The argument N may be a template parameter. This attribute should be +ignored for the FPGA emulator device. + +``[[intelfpga::scheduler_target_fmax_mhz(N)]]`` +Valid values of N are integers in the range [0, 1048576]. The upper limit, +although too high to be a realistic value for frequency, is chosen to be future +proof. The FPGA backend emits a diagnostic message if the passed value is +unachievable by the device. + +This attribute enables communication of the desired maximum frequency of the +device operation, guiding the FPGA backend to insert the appropriate number of +registers to break-up the combinational logic circuit, and thereby controlling +the length of the longest combinational path. + }]; +} + def SYCLIntelNoGlobalWorkOffsetAttrDocs : Documentation { let Category = DocCatFunction; let Heading = "no_global_work_offset (IntelFPGA)"; diff --git a/clang/include/clang/Basic/AttributeCommonInfo.h b/clang/include/clang/Basic/AttributeCommonInfo.h index 2fbaa8dfc22fa..1884813f2524a 100644 --- a/clang/include/clang/Basic/AttributeCommonInfo.h +++ b/clang/include/clang/Basic/AttributeCommonInfo.h @@ -162,6 +162,7 @@ class AttributeCommonInfo { (ParsedAttr == AT_ReqdWorkGroupSize && isCXX11Attribute()) || (ParsedAttr == AT_IntelReqdSubGroupSize && isCXX11Attribute()) || ParsedAttr == AT_SYCLIntelNumSimdWorkItems || + ParsedAttr == AT_SYCLIntelSchedulerTargetFmaxMhz || ParsedAttr == AT_SYCLIntelMaxWorkGroupSize || ParsedAttr == AT_SYCLIntelMaxGlobalWorkDim || ParsedAttr == AT_SYCLIntelNoGlobalWorkOffset) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c565bc055c88e..8490e7b97eb40 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10045,6 +10045,11 @@ class Sema final { bool checkAllowedSYCLInitializer(VarDecl *VD, bool CheckValueDependent = false); + // Adds a scheduler_target_fmax_mhz attribute to a particular declaration. + void addSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D, + const AttributeCommonInfo &CI, + Expr *E); + //===--------------------------------------------------------------------===// // C++ Coroutines TS // diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index bdd1a43ccf7cb..542927b5a7ea9 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -639,6 +639,17 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD, llvm::MDNode::get(Context, AttrMDArgs)); } + if (const SYCLIntelSchedulerTargetFmaxMhzAttr *A = + FD->getAttr()) { + Optional ArgVal = + A->getValue()->getIntegerConstantExpr(FD->getASTContext()); + assert(ArgVal.hasValue() && "Not an integer constant expression"); + llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get( + Builder.getInt32(ArgVal->getSExtValue()))}; + Fn->setMetadata("scheduler_target_fmax_mhz", + llvm::MDNode::get(Context, AttrMDArgs)); + } + if (const SYCLIntelMaxWorkGroupSizeAttr *A = FD->getAttr()) { llvm::Metadata *AttrMDArgs[] = { diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 8925c3ff895d0..02e98bb37cb9d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3006,6 +3006,38 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D, E); } +// Add scheduler_target_fmax_mhz +void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr( + Decl *D, const AttributeCommonInfo &Attr, Expr *E) { + assert(E && "Attribute must have an argument."); + + SYCLIntelSchedulerTargetFmaxMhzAttr TmpAttr(Context, Attr, E); + if (!E->isValueDependent()) { + ExprResult ResultExpr; + if (checkRangedIntegralArgument( + E, &TmpAttr, ResultExpr)) + return; + E = ResultExpr.get(); + } + + D->addAttr(::new (Context) + SYCLIntelSchedulerTargetFmaxMhzAttr(Context, Attr, E)); +} + +// Handle scheduler_target_fmax_mhz +static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D, + const ParsedAttr &AL) { + if (D->isInvalidDecl()) + return; + + Expr *E = AL.getArgAsExpr(0); + + if (D->getAttr()) + S.Diag(AL.getLoc(), diag::warn_duplicate_attribute) << AL; + + S.addSYCLIntelSchedulerTargetFmaxMhzAttr(D, AL, E); +} + // Handles max_global_work_dim. static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D, const ParsedAttr &Attr) { @@ -7893,6 +7925,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLIntelNumSimdWorkItems: handleNumSimdWorkItemsAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLIntelSchedulerTargetFmaxMhz: + handleSchedulerTargetFmaxMhzAttr(S, D, AL); + break; case ParsedAttr::AT_SYCLIntelMaxGlobalWorkDim: handleMaxGlobalWorkDimAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c52c4c229ec10..6e25ae7a6a975 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -531,6 +531,9 @@ class MarkDeviceFunction : public RecursiveASTVisitor { if (auto *A = FD->getAttr()) Attrs.insert(A); + if (auto *A = FD->getAttr()) + Attrs.insert(A); + if (auto *A = FD->getAttr()) Attrs.insert(A); @@ -3166,6 +3169,7 @@ void Sema::MarkDevice(void) { } case attr::Kind::SYCLIntelKernelArgsRestrict: case attr::Kind::SYCLIntelNumSimdWorkItems: + case attr::Kind::SYCLIntelSchedulerTargetFmaxMhz: case attr::Kind::SYCLIntelMaxGlobalWorkDim: case attr::Kind::SYCLIntelNoGlobalWorkOffset: case attr::Kind::SYCLSimd: { diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 26dd65fa016f8..9ca764418a3f1 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -737,6 +737,12 @@ void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs, *this, TemplateArgs, SYCLIntelNumSimdWorkItems, New); continue; } + if (const auto *SYCLIntelSchedulerTargetFmaxMhz = + dyn_cast(TmplAttr)) { + instantiateIntelSYCLFunctionAttr( + *this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New); + continue; + } // Existing DLL attribute on the instantiation takes precedence. if (TmplAttr->getKind() == attr::DLLExport || TmplAttr->getKind() == attr::DLLImport) { diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp new file mode 100644 index 0000000000000..d90b09e94bb29 --- /dev/null +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -0,0 +1,25 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s + +#include "Inputs/sycl.hpp" +[[intelfpga::scheduler_target_fmax_mhz(5)]] void +func() {} + +template +[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {} + +int main() { + cl::sycl::kernel_single_task( + []() [[intelfpga::scheduler_target_fmax_mhz(2)]]{}); + + cl::sycl::kernel_single_task( + []() { func(); }); + + cl::sycl::kernel_single_task( + []() { zoo<75>(); }); +} +// CHECK: define spir_kernel void @{{.*}}test_kernel1() {{.*}} !scheduler_target_fmax_mhz ![[PARAM1:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}test_kernel2() {{.*}} !scheduler_target_fmax_mhz ![[PARAM2:[0-9]+]] +// CHECK: define spir_kernel void @{{.*}}test_kernel3() {{.*}} !scheduler_target_fmax_mhz ![[PARAM3:[0-9]+]] +// CHECK: ![[PARAM1]] = !{i32 2} +// CHECK: ![[PARAM2]] = !{i32 5} +// CHECK: ![[PARAM3]] = !{i32 75} diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp new file mode 100644 index 0000000000000..d6480849c796d --- /dev/null +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp @@ -0,0 +1,45 @@ +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s + +#include "Inputs/sycl.hpp" +[[intelfpga::scheduler_target_fmax_mhz(2)]] void +func() {} + +template +[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {} + +int main() { + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 'void ()' + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 5 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5 + cl::sycl::kernel_single_task( + []() [[intelfpga::scheduler_target_fmax_mhz(5)]]{}); + + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 'void ()' + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 + cl::sycl::kernel_single_task( + []() { func(); }); + + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 'void ()' + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int' + // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75 + cl::sycl::kernel_single_task( + []() { zoo<75>(); }); + + [[intelfpga::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}} + + cl::sycl::kernel_single_task( + []() [[intelfpga::scheduler_target_fmax_mhz(1048577)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}} + + cl::sycl::kernel_single_task( + []() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires integer constant between 0 and 1048576 inclusive}} + + cl::sycl::kernel_single_task( + []() [[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}} +}