[SYCL] Add clang support for FPGA kernel attribute scheduler_target_fmax_mhz

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">,

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 therby controlling
+the length of the longest combinational path.
+  }];
+}
+
 def SYCLIntelNoGlobalWorkOffsetAttrDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "no_global_work_offset (IntelFPGA)";

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)

clang/include/clang/Sema/Sema.h

@@ -10045,6 +10045,12 @@ class Sema final {
   bool checkAllowedSYCLInitializer(VarDecl *VD,
                                    bool CheckValueDependent = false);
 
+  // Adds an scheduler_target_fmax_mhz intel_reqd_sub_group_size attribute to a
+  // particular declaration.
+  void addSYCLIntelSchedulerTargetFmaxMhzAttr(Decl *D,
+                                              const AttributeCommonInfo &CI,
+                                              Expr *E);
+
   //===--------------------------------------------------------------------===//
   // C++ Coroutines TS
   //

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<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
+    Optional<llvm::APSInt> 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<SYCLIntelMaxWorkGroupSizeAttr>()) {
     llvm::Metadata *AttrMDArgs[] = {

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3006,6 +3006,39 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
                                                                      E);
 }
 
+// Add scheduler_target_fmax_mhz
+void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr(
+    Decl *D, const AttributeCommonInfo &Attr, Expr *E) {
+  if (!E)
+    return;
+
+  SYCLIntelSchedulerTargetFmaxMhzAttr TmpAttr(Context, Attr, E);
+  if (!E->isValueDependent()) {
+    ExprResult ResultExpr;
+    if (checkRangedIntegralArgument<SYCLIntelSchedulerTargetFmaxMhzAttr>(
+            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<SYCLIntelSchedulerTargetFmaxMhzAttr>())
+    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 +7926,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;
@@ -8314,6 +8350,10 @@ void Sema::ProcessDeclAttributeList(Scope *S, Decl *D,
         Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
         D->setInvalidDecl();
       }
+    } else if (const auto *A =
+                   D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>()) {
+      Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
+      D->setInvalidDecl();
     } else if (!D->hasAttr<CUDAGlobalAttr>()) {
       if (const auto *A = D->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
         Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)

clang/lib/Sema/SemaSYCL.cpp

@@ -531,6 +531,9 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
       if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
         Attrs.insert(A);
 
+      if (auto *A = FD->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
+        Attrs.insert(A);
+
       if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
         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: {

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<SYCLIntelSchedulerTargetFmaxMhzAttr>(TmplAttr)) {
+      instantiateIntelSYCLFunctionAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>(
+          *this, TemplateArgs, SYCLIntelSchedulerTargetFmaxMhz, New);
+      continue;
+    }
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {

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 <int N>
+[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {}
+
+int main() {
+  cl::sycl::kernel_single_task<class test_kernel1>(
+      []() [[intelfpga::scheduler_target_fmax_mhz(2)]]{});
+
+  cl::sycl::kernel_single_task<class test_kernel2>(
+      []() { func(); });
+
+  cl::sycl::kernel_single_task<class test_kernel3>(
+      []() { 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}

clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

@@ -0,0 +1,52 @@
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
+
+#include "Inputs/sycl.hpp"
+#ifndef TRIGGER_ERROR
+[[intelfpga::scheduler_target_fmax_mhz(2)]] // expected-no-diagnostics
+void
+func() {}
+
+template <int N>
+[[intelfpga::scheduler_target_fmax_mhz(N)]] void zoo() {}
+#endif // TRIGGER_ERROR
+
+int main() {
+#ifndef TRIGGER_ERROR
+  // 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<class test_kernel1>(
+      []() [[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<class test_kernel2>(
+      []() { 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<class test_kernel3>(
+      []() { zoo<75>(); });
+#else
+  [[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<class test_kernel4>(
+      []() [[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<class test_kernel5>(
+      []() [[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<class test_kernel6>(
+      []() [[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}}
+#endif // TRIGGER_ERROR
+}