Complete draft patch and add a test

Author: schittir

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/lib/Sema/SemaDeclAttr.cpp

@@ -3005,6 +3005,30 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
   S.addIntelSYCLSingleArgFunctionAttr<SYCLIntelNumSimdWorkItemsAttr>(D, Attr,
                                                                      E);
 }
+// Handles scheduler_target_fmax_mhz
+static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
+                                       const ParsedAttr &Attr) {
+  if (D->isInvalidDecl())
+    return;
+
+  uint32_t TargetFmaxMhz = 0;
+  const Expr *E = Attr.getArgAsExpr(0);
+  if (!checkUInt32Argument(S, Attr, E, TargetFmaxMhz, 0,
+                           /*StrictlyUnsigned=*/true))
+    return;
+
+  if (TargetFmaxMhz == 0) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
+        << Attr << E->getSourceRange();
+    return;
+  }
+
+  if (D->getAttr<SYCLIntelSchedulerTargetFmaxMhzAttr>())
+    S.Diag(Attr.getLoc(), diag::warn_duplicate_attribute) << Attr;
+
+  D->addAttr(::new (S.Context) SYCLIntelSchedulerTargetFmaxMhzAttr(
+      S.Context, Attr, TargetFmaxMhz));
+}
 
 // Handles max_global_work_dim.
 static void handleMaxGlobalWorkDimAttr(Sema &S, Decl *D,
@@ -7893,6 +7917,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;

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/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

@@ -0,0 +1,36 @@
+// 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() {}
+#endif // TRIGGER_ERROR
+
+int main() {
+#ifndef TRIGGER_ERROR
+  // CHECK-LABEL:  -FunctionDecl {{.*}}test_kernel1 'void ()'
+  // CHECK:        SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} 5
+  cl::sycl::kernel_single_task<class test_kernel1>(
+      []() [[intelfpga::scheduler_target_fmax_mhz(5)]]{});
+
+  // CHECK-LABEL:  `-FunctionDecl {{.*}}test_kernel2 'void ()'
+  // CHECK:  -SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} 2
+  cl::sycl::kernel_single_task<class test_kernel2>(
+      []() { func(); });
+
+#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_kernel3>(
+      []() [[intelfpga::scheduler_target_fmax_mhz(0)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute must be greater than 0}}
+
+  cl::sycl::kernel_single_task<class test_kernel4>(
+      []() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
+
+  cl::sycl::kernel_single_task<class test_kernel5>(
+      []() [[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
+}