Expand attribute defintion and add upper limit for attribute parameter

Author: schittir

clang/include/clang/Basic/AttrDocs.td

@@ -2205,6 +2205,17 @@ 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.
   }];
 }
 

clang/lib/Sema/SemaDeclAttr.cpp

@@ -3017,9 +3017,12 @@ static void handleSchedulerTargetFmaxMhzAttr(Sema &S, Decl *D,
                            /*StrictlyUnsigned=*/true))
     return;
 
-  if (TargetFmaxMhz == 0) {
-    S.Diag(Attr.getLoc(), diag::err_attribute_argument_is_zero)
-        << Attr << E->getSourceRange();
+  uint32_t UpperLimit = 1048576;
+  if (TargetFmaxMhz > UpperLimit) {
+    // Allow frequency to be a "large enough" positive value, zero included
+    // Let FPGA backend handle unachievable frequency value
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_out_of_range)
+        << Attr << E->getSourceRange() << 0 << UpperLimit;
     return;
   }
 

clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

@@ -25,7 +25,7 @@ int main() {
   [[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}}
+      []() [[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_kernel4>(
       []() [[intelfpga::scheduler_target_fmax_mhz(-4)]]{}); // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}