Add diagnostics for conflicting attributes from transitive calls

Author: schittir

clang/lib/Sema/SemaSYCL.cpp

@@ -2027,6 +2027,33 @@ 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);
+          }
+          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);
           }
@@ -2035,7 +2062,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

@@ -1,13 +1,39 @@
+// 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
 
-[[intelfpga::no_global_work_offset]] void not_direct() {}
+#ifndef TRIGGER_ERROR
+[[intelfpga::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics
 
-void func() { not_direct(); }
+[[cl::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();
-  func();
+#ifndef TRIGGER_ERROR
+  func_one();
+#else
+  func_two();
+  func_three();
+#endif
 }
 
 template <typename Name, typename Type>
@@ -16,8 +42,13 @@ void parallel_for(Type lambda) {
 }
 
 void invoke_foo2() {
+#ifndef TRIGGER_ERROR
   // CHECK-LABEL:  FunctionDecl {{.*}} invoke_foo2 'void ()'
   // CHECK:        `-FunctionDecl {{.*}}KernelName 'void ()'
-  // CHECK:        -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Enabled
+  // 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
 }