[SYCL] Propagate attributes from transitive calls to kernel

clang/lib/Sema/SemaSYCL.cpp

@@ -514,52 +514,25 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
 
       if (auto *A = FD->getAttr<IntelReqdSubGroupSizeAttr>())
         Attrs.insert(A);
+
       if (auto *A = FD->getAttr<ReqdWorkGroupSizeAttr>())
         Attrs.insert(A);
-      // Allow the following kernel attributes only on lambda functions and
-      // function objects that are called directly from a kernel (i.e. the one
-      // passed to the parallel_for function). For all other cases,
-      // emit a warning and ignore.
-      if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelKernelArgsRestrictAttr>();
-        }
-      }
-      if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelNumSimdWorkItemsAttr>();
-        }
-      }
-      if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelMaxWorkGroupSizeAttr>();
-        }
-      }
-      if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelMaxGlobalWorkDimAttr>();
-        }
-      }
-      if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>()) {
-        if (ParentFD == SYCLKernel) {
-          Attrs.insert(A);
-        } else {
-          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
-          FD->dropAttr<SYCLIntelNoGlobalWorkOffsetAttr>();
-        }
-      }
+
+      if (auto *A = FD->getAttr<SYCLIntelKernelArgsRestrictAttr>())
+        Attrs.insert(A);
+
+      if (auto *A = FD->getAttr<SYCLIntelNumSimdWorkItemsAttr>())
+        Attrs.insert(A);
+
+      if (auto *A = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>())
+        Attrs.insert(A);
+
+      if (auto *A = FD->getAttr<SYCLIntelMaxGlobalWorkDimAttr>())
+        Attrs.insert(A);
+
+      if (auto *A = FD->getAttr<SYCLIntelNoGlobalWorkOffsetAttr>())
+        Attrs.insert(A);
+
       if (auto *A = FD->getAttr<SYCLSimdAttr>())
         Attrs.insert(A);
       // Propagate the explicit SIMD attribute through call graph - it is used
@@ -2054,6 +2027,38 @@ void Sema::MarkDevice(void) {
               Diag(Attr->getLocation(), diag::note_conflicting_attribute);
               SYCLKernel->setInvalidDecl();
             }
+          } else if (auto *Existing =
+                         SYCLKernel->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
+            if (Existing->getXDim() < Attr->getXDim() ||
+                Existing->getYDim() < Attr->getYDim() ||
+                Existing->getZDim() < Attr->getZDim()) {
+              Diag(SYCLKernel->getLocation(),
+                   diag::err_conflicting_sycl_kernel_attributes);
+              Diag(Existing->getLocation(), diag::note_conflicting_attribute);
+              Diag(Attr->getLocation(), diag::note_conflicting_attribute);
+              SYCLKernel->setInvalidDecl();
+            } else {
+              SYCLKernel->addAttr(A);
+            }
+          } else {
+            SYCLKernel->addAttr(A);
+          }
+          break;
+        }
+        case attr::Kind::SYCLIntelMaxWorkGroupSize: {
+          auto *Attr = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
+          if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {
+            if (Existing->getXDim() > Attr->getXDim() ||
+                Existing->getYDim() > Attr->getYDim() ||
+                Existing->getZDim() > Attr->getZDim()) {
+              Diag(SYCLKernel->getLocation(),
+                   diag::err_conflicting_sycl_kernel_attributes);
+              Diag(Existing->getLocation(), diag::note_conflicting_attribute);
+              Diag(Attr->getLocation(), diag::note_conflicting_attribute);
+              SYCLKernel->setInvalidDecl();
+            } else {
+              SYCLKernel->addAttr(A);
+            }
           } else {
             SYCLKernel->addAttr(A);
           }
@@ -2062,7 +2067,6 @@ void Sema::MarkDevice(void) {
         case attr::Kind::SYCLIntelKernelArgsRestrict:
         case attr::Kind::SYCLIntelNumSimdWorkItems:
         case attr::Kind::SYCLIntelMaxGlobalWorkDim:
-        case attr::Kind::SYCLIntelMaxWorkGroupSize:
         case attr::Kind::SYCLIntelNoGlobalWorkOffset:
         case attr::Kind::SYCLSimd: {
           if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&

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

@@ -0,0 +1,54 @@
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -verify
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
+
+#ifndef TRIGGER_ERROR
+[[intelfpga::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics
+
+[[intel::reqd_sub_group_size(1)]] void func_one() {
+  not_direct_one();
+}
+
+#else
+[[cl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note {{conflicting attribute is here}}
+
+[[intelfpga::max_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
+void
+func_two() {
+  not_direct_two();
+}
+
+[[cl::reqd_work_group_size(4, 4, 4)]] // expected-note 2 {{conflicting attribute is here}}
+void
+func_three() {
+  not_direct_two();
+}
+#endif
+
+template <typename Name, typename Type>
+[[clang::sycl_kernel]] void __my_kernel__(Type bar) {
+  bar();
+#ifndef TRIGGER_ERROR
+  func_one();
+#else
+  func_two();
+  func_three();
+#endif
+}
+
+template <typename Name, typename Type>
+void parallel_for(Type lambda) {
+  __my_kernel__<Name>(lambda);
+}
+
+void invoke_foo2() {
+#ifndef TRIGGER_ERROR
+  // CHECK-LABEL:  FunctionDecl {{.*}} invoke_foo2 'void ()'
+  // CHECK:        `-FunctionDecl {{.*}}KernelName 'void ()'
+  // CHECK:        -IntelReqdSubGroupSizeAttr {{.*}}
+  // CHECK:        `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
+  parallel_for<class KernelName>([]() {});
+#else
+  parallel_for<class KernelName>([]() {}); // expected-error 2 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
+#endif
+}

clang/test/SemaSYCL/intel-max-global-work-dim.cpp

@@ -20,8 +20,7 @@ void foo() {
 
 #else // __SYCL_DEVICE_ONLY__
 
-[[intelfpga::max_global_work_dim(2)]] // expected-warning{{'max_global_work_dim' attribute ignored}}
-void func_ignore() {}
+[[intelfpga::max_global_work_dim(2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
   [[intelfpga::max_global_work_dim(1)]]
@@ -68,9 +67,9 @@ int main() {
       []() [[intelfpga::max_global_work_dim(2)]] {});
 
   // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
-  // CHECK-NOT:   SYCLIntelMaxGlobalWorkDimAttr {{.*}}
+  // CHECK:       SYCLIntelMaxGlobalWorkDimAttr {{.*}}
   kernel<class test_kernel3>(
-      []() {func_ignore();});
+      []() { func_do_not_ignore(); });
 
   kernel<class test_kernel4>(
       TRIFuncObjGood1());

clang/test/SemaSYCL/intel-max-work-group-size.cpp

@@ -20,8 +20,7 @@ void foo() {
 
 #else // __SYCL_DEVICE_ONLY__
 
-[[intelfpga::max_work_group_size(2, 2, 2)]] // expected-warning{{'max_work_group_size' attribute ignored}}
-void func_ignore() {}
+[[intelfpga::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
   [[intelfpga::max_work_group_size(4, 4, 4)]]
@@ -53,9 +52,9 @@ int main() {
       []() [[intelfpga::max_work_group_size(8, 8, 8)]] {});
 
   // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
-  // CHECK-NOT:   SYCLIntelMaxWorkGroupSizeAttr {{.*}}
+  // CHECK:       SYCLIntelMaxWorkGroupSizeAttr {{.*}}
   kernel<class test_kernel3>(
-      []() {func_ignore();});
+      []() { func_do_not_ignore(); });
 
 #ifdef TRIGGER_ERROR
   [[intelfpga::max_work_group_size(1, 1, 1)]] int Var = 0; // expected-error{{'max_work_group_size' attribute only applies to functions}}

clang/test/SemaSYCL/intel-restrict.cpp

@@ -1,8 +1,7 @@
 // RUN: %clang_cc1 %s -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DCHECKDIAG -verify
 // RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 | FileCheck %s
 
-[[intel::kernel_args_restrict]] // expected-warning{{'kernel_args_restrict' attribute ignored}}
-void func_ignore() {}
+[[intel::kernel_args_restrict]] void func_do_not_ignore() {}
 
 struct FuncObj {
   [[intel::kernel_args_restrict]]
@@ -29,7 +28,7 @@ int main() {
       []() [[intel::kernel_args_restrict]] {});
 
   // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
-  // CHECK-NOT:   SYCLIntelKernelArgsRestrictAttr
+  // CHECK:       SYCLIntelKernelArgsRestrictAttr
   kernel<class test_kernel3>(
-      []() {func_ignore();});
+      []() { func_do_not_ignore(); });
 }

clang/test/SemaSYCL/num_simd_work_items.cpp

@@ -20,8 +20,7 @@ void foo() {
 
 #else // __SYCL_DEVICE_ONLY__
 
-[[intelfpga::num_simd_work_items(2)]] // expected-warning{{'num_simd_work_items' attribute ignored}}
-void func_ignore() {}
+[[intelfpga::num_simd_work_items(2)]] void func_do_not_ignore() {}
 
 struct FuncObj {
   [[intelfpga::num_simd_work_items(42)]]
@@ -45,9 +44,9 @@ int main() {
       []() [[intelfpga::num_simd_work_items(8)]] {});
 
   // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
-  // CHECK-NOT:   SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
+  // CHECK:       SYCLIntelNumSimdWorkItemsAttr {{.*}} 2
   kernel<class test_kernel3>(
-      []() {func_ignore();});
+      []() { func_do_not_ignore(); });
 
 #ifdef TRIGGER_ERROR
   [[intelfpga::num_simd_work_items(0)]] int Var = 0; // expected-error{{'num_simd_work_items' attribute only applies to functions}}