[SYCL][FPGA] [WIP]Do not propagate the attributes from device functions to a kernel

clang/include/clang/Basic/AttrDocs.td

@@ -2418,7 +2418,8 @@ attribute was applied. This effect is equivalent to annotating restrict on
 **all** kernel pointer arguments in an OpenCL or SPIR-V kernel.
 
 If ``intel::kernel_args_restrict`` is applied to a function called from a device
-kernel, the attribute is not ignored and it is propagated to the kernel.
+kernel, the attribute is ignored and it is not propagated to the kernel to
+match with new SYCL 2020 spec.
 
 The attribute forms an unchecked assertion, in that implementations
 do not need to check/confirm the pre-condition in any way. If a user applies
@@ -2711,6 +2712,10 @@ 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.
 
+If ``intel::scheduler_target_fmax_mhz`` is applied to a function called from a
+device kernel, the attribute is ignored and it is not propagated to the kernel
+to match with new SYCL 2020 spec.
+
 .. code-block:: c++
 
   [[intel::scheduler_target_fmax_mhz(4)]] void foo() {}
@@ -2741,6 +2746,10 @@ function object). If 1, compiler doesn't use the global work offset values for
 the device function. Valid values are 0 and 1. If used without argument, value
 of 1 is set implicitly.
 
+If ``intel::no_global_work_offset`` is applied to a function called from a
+device kernel, the attribute is ignored and it is not propagated to the kernel
+to match with new SYCL 2020 spec.
+
 .. code-block:: c++
 
   [[intel::no_global_work_offset]]

clang/lib/Sema/SemaSYCL.cpp

@@ -346,11 +346,9 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
   llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
     // FIXME: Make this list self-adapt as new SYCL attributes are added.
     return isa<IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr,
-               ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr,
-               SYCLIntelNumSimdWorkItemsAttr,
-               SYCLIntelSchedulerTargetFmaxMhzAttr,
+               ReqdWorkGroupSizeAttr, SYCLIntelNumSimdWorkItemsAttr,
                SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr,
-               SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A);
+               SYCLSimdAttr>(A);
   });
 
   // Allow the kernel attribute "use_stall_enable_clusters" only on lambda
@@ -372,6 +370,9 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD,
     llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) {
       return isa<SYCLIntelLoopFuseAttr, SYCLIntelFPGAMaxConcurrencyAttr,
                  SYCLIntelFPGADisableLoopPipeliningAttr,
+                 SYCLIntelKernelArgsRestrictAttr,
+                 SYCLIntelNoGlobalWorkOffsetAttr,
+                 SYCLIntelSchedulerTargetFmaxMhzAttr,
                  SYCLIntelFPGAInitiationIntervalAttr>(A);
     });
   }

clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
 
 #include "sycl.hpp"
 
@@ -16,8 +16,7 @@ class Functor {
   [[intel::no_global_work_offset(SIZE)]] void operator()() const {}
 };
 
-template <int N>
-[[intel::no_global_work_offset(N)]] void func() {}
+[[intel::no_global_work_offset(1)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
@@ -33,9 +32,9 @@ int main() {
     Functor<1> f;
     h.single_task<class kernel_name4>(f);
 
-    h.single_task<class kernel_name5>([]() {
-      func<1>();
-    });
+    // Test attribute is not propagated.
+    h.single_task<class kernel_name5>(
+        []() { func(); });
   });
   return 0;
 }
@@ -44,6 +43,6 @@ int main() {
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]]
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
-// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} ![[NUM5]]
 // CHECK-NOT: ![[NUM4]]  = !{i32 0}
 // CHECK: ![[NUM5]] = !{}

clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp

@@ -1,25 +1,49 @@
-// RUN: %clang_cc1 -fsycl-is-device -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s
 
-#include "Inputs/sycl.hpp"
-[[intel::scheduler_target_fmax_mhz(5)]] void
-func() {}
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+class Foo {
+public:
+  [[intel::scheduler_target_fmax_mhz(5)]] void operator()() const {}
+};
 
 template <int N>
-[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {}
+class Functor {
+public:
+  [[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {}
+};
+
+[[intel::scheduler_target_fmax_mhz(2)]] void zoo() {}
 
 int main() {
-  cl::sycl::kernel_single_task<class test_kernel1>(
-      []() [[intel::scheduler_target_fmax_mhz(2)]]{});
+  q.submit([&](handler &h) {
+    // Test attribute argument size.
+    Foo boo;
+    h.single_task<class kernel_name1>(boo);
 
-  cl::sycl::kernel_single_task<class test_kernel2>(
-      []() { func(); });
+    // Test attribute is applied on lambda.
+    h.single_task<class kernel_name2>(
+        []() [[intel::scheduler_target_fmax_mhz(42)]]{});
 
-  cl::sycl::kernel_single_task<class test_kernel3>(
-      []() { zoo<75>(); });
+    // Test template argument.
+    Functor<75> f;
+    h.single_task<class kernel_name3>(f);
+
+    // Test attribute is not propagated.
+    h.single_task<class kernel_name4>(
+        []() { zoo(); });
+  });
+  return 0;
 }
-// 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}
+
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]]
+// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} ![[NUM0:[0-9]+]]
+// CHECK: ![[NUM0]] = !{}
+// CHECK: ![[NUM5]] = !{i32 5}
+// CHECK: ![[NUM42]] = !{i32 42}
+// CHECK: ![[NUM75]] = !{i32 75}

clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp

@@ -46,7 +46,7 @@ void invoke_foo2() {
   // CHECK-LABEL:  FunctionDecl {{.*}} invoke_foo2 'void ()'
   // CHECK:        `-FunctionDecl {{.*}}KernelName 'void ()'
   // CHECK:        -IntelReqdSubGroupSizeAttr {{.*}}
-  // CHECK:        `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
+  // CHECK-NOT:    `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
   parallel_for<class KernelName>([]() {});
 #else
   parallel_for<class KernelName>([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}

clang/test/SemaSYCL/intel-restrict.cpp

@@ -27,7 +27,7 @@ int main() {
       []() [[intel::kernel_args_restrict]] {});
 
   // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
-  // CHECK:       SYCLIntelKernelArgsRestrictAttr
+  // CHECK-NOT:   SYCLIntelKernelArgsRestrictAttr
   kernel<class test_kernel3>(
       []() { func_do_not_ignore(); });
 }

clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp

@@ -53,10 +53,7 @@ int main() {
 #ifndef TRIGGER_ERROR
     // CHECK-LABEL:  FunctionDecl {{.*}} main 'int ()'
     // CHECK:  `-FunctionDecl {{.*}}test_kernel1 'void ()'
-    // CHECK:  -SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
-    // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
-    // CHECK-NEXT:  value: Int 1
-    // CHECK-NEXT:  IntegerLiteral{{.*}}1{{$}}
+    // CHECK-NOT:   SYCLIntelNoGlobalWorkOffsetAttr {{.*}}
     // CHECK:  -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited
     // CHECK-NEXT:  ConstantExpr {{.*}} 'int'
     // CHECK-NEXT:  value: Int 4

clang/test/SemaSYCL/scheduler_target_fmax_mhz.cpp

@@ -1,60 +1,77 @@
-// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -verify %s
 
-#include "Inputs/sycl.hpp"
+// Test that checks scheduler_target_fmax_mhz attribute support on Function.
+
+// Test for deprecated spelling of scheduler_target_fmax_mhz attribute.
 // expected-warning@+2 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
 // expected-note@+1 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
-[[intelfpga::scheduler_target_fmax_mhz(2)]] void
-func() {}
+[[intelfpga::scheduler_target_fmax_mhz(2)]] void deprecate() {}
+
+// Tests for incorrect argument values for Intel FPGA scheduler_target_fmax_mhz function attribute.
+[[intel::scheduler_target_fmax_mhz(0)]] int Var = 0; // expected-error{{'scheduler_target_fmax_mhz' attribute only applies to functions}}
+
+[[intel::scheduler_target_fmax_mhz(1048577)]] void correct() {}  // OK
+
+[[intel::scheduler_target_fmax_mhz("foo")]] void func() {} // expected-error{{integral constant expression must have integral or unscoped enumeration type, not 'const char [4]'}}
+
+[[intel::scheduler_target_fmax_mhz(-1)]] void func1() {} // expected-error{{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
+
+[[intel::scheduler_target_fmax_mhz(0, 1)]] void func2() {} // expected-error{{'scheduler_target_fmax_mhz' attribute takes one argument}}
+
+// Tests for Intel FPGA scheduler_target_fmax_mhz function attribute duplication.
+// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
+[[intel::scheduler_target_fmax_mhz(2)]]
+[[intel::scheduler_target_fmax_mhz(2)]] void func3() {}
 
 // No diagnostic is emitted because the arguments match.
-[[intel::scheduler_target_fmax_mhz(12)]] void bar();
-[[intel::scheduler_target_fmax_mhz(12)]] void bar() {} // OK
+[[intel::scheduler_target_fmax_mhz(4)]] void func4();
+[[intel::scheduler_target_fmax_mhz(4)]] void func4(); // OK
 
 // Diagnostic is emitted because the arguments mismatch.
-[[intel::scheduler_target_fmax_mhz(12)]] void baz();  // expected-note {{previous attribute is here}}
-[[intel::scheduler_target_fmax_mhz(100)]] void baz(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
+[[intel::scheduler_target_fmax_mhz(2)]]                  // expected-note {{previous attribute is here}}
+[[intel::scheduler_target_fmax_mhz(4)]] void func5() {}  // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
+
+[[intel::scheduler_target_fmax_mhz(1)]] void func6(); // expected-note {{previous attribute is here}}
+[[intel::scheduler_target_fmax_mhz(3)]] void func6(); // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
 
+// Tests that check template parameter support for Intel FPGA scheduler_target_fmax_mhz function attributes.
 template <int N>
-[[intel::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
-  // expected-warning@+3 {{attribute 'intelfpga::scheduler_target_fmax_mhz' is deprecated}}
-  // expected-note@+2 {{did you mean to use 'intel::scheduler_target_fmax_mhz' instead?}}
-  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:   ConstantExpr {{.*}} 'int'
-  // CHECK-NEXT:   value: Int 75
-  // 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>(); });
-
-  [[intel::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>(
-      []() [[intel::scheduler_target_fmax_mhz(1048577)]]{}); // OK
-
-  cl::sycl::kernel_single_task<class test_kernel5>(
-      []() [[intel::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_kernel6>(
-      []() [[intel::scheduler_target_fmax_mhz(1),      // expected-note {{previous attribute is here}}
-             intel::scheduler_target_fmax_mhz(2)]]{}); // expected-warning{{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
+[[intel::scheduler_target_fmax_mhz(N)]] void func7(); // expected-error {{'scheduler_target_fmax_mhz' attribute requires a non-negative integral compile time constant expression}}
+
+template <int size>
+[[intel::scheduler_target_fmax_mhz(10)]] void func8();     // expected-note {{previous attribute is here}}
+template <int size>
+[[intel::scheduler_target_fmax_mhz(size)]] void func8() {} // expected-warning {{attribute 'scheduler_target_fmax_mhz' is already applied with different arguments}}
+
+void checkTemplates() {
+  func7<4>();  // OK
+  func7<-1>(); // expected-note {{in instantiation of function template specialization 'func7<-1>' requested here}}
+  func7<0>();  // OK
+  func8<20>(); // expected-note {{in instantiation of function template specialization 'func8<20>' requested here}}
+}
+
+// Test that checks expression is not a constant expression.
+// expected-note@+1{{declared here}}
+int baz();
+// expected-error@+2{{expression is not an integral constant expression}}
+// expected-note@+1{{non-constexpr function 'baz' cannot be used in a constant expression}}
+[[intel::scheduler_target_fmax_mhz(baz() + 1)]] void func9();
+
+// Test that checks expression is a constant expression.
+constexpr int bar() { return 0; }
+[[intel::scheduler_target_fmax_mhz(bar() + 2)]] void func10(); // OK
+
+// Test that checks wrong function template instantiation and ensures that the type
+// is checked properly when instantiating from the template definition.
+template <typename Ty>
+// expected-error@+2 {{integral constant expression must have integral or unscoped enumeration type, not 'S'}}
+// expected-error@+1 {{integral constant expression must have integral or unscoped enumeration type, not 'float'}}
+[[intel::scheduler_target_fmax_mhz(Ty{})]] void func11() {}
+
+struct S {};
+void test() {
+  //expected-note@+1{{in instantiation of function template specialization 'func11<S>' requested here}}
+  func11<S>();
+  //expected-note@+1{{in instantiation of function template specialization 'func11<float>' requested here}}
+  func11<float>();
 }