[SYCL] Propagate attributes from redeclarations to SYCL kernel (#2063)

- Fix bug that propagates attributes from only canoncial declaration
- Handle conflicting attributes
- Add a test

Author: schittir

clang/lib/Sema/SemaDecl.cpp

@@ -3212,6 +3212,22 @@ static void adjustDeclContextForDeclaratorDecl(DeclaratorDecl *NewD,
     FixSemaDC(VD->getDescribedVarTemplate());
 }
 
+template <typename AttributeType>
+static void checkDimensionsAndSetDiagnostics(Sema &S, FunctionDecl *New,
+                                             FunctionDecl *Old) {
+  AttributeType *NewDeclAttr = New->getAttr<AttributeType>();
+  AttributeType *OldDeclAttr = Old->getAttr<AttributeType>();
+  if ((NewDeclAttr->getXDim() != OldDeclAttr->getXDim()) ||
+      (NewDeclAttr->getYDim() != OldDeclAttr->getYDim()) ||
+      (NewDeclAttr->getZDim() != OldDeclAttr->getZDim())) {
+    S.Diag(New->getLocation(), diag::err_conflicting_sycl_function_attributes)
+        << OldDeclAttr << NewDeclAttr;
+    S.Diag(New->getLocation(), diag::warn_duplicate_attribute) << OldDeclAttr;
+    S.Diag(OldDeclAttr->getLocation(), diag::note_conflicting_attribute);
+    S.Diag(NewDeclAttr->getLocation(), diag::note_conflicting_attribute);
+  }
+}
+
 /// MergeFunctionDecl - We just parsed a function 'New' from
 /// declarator D which has the same name and scope as a previous
 /// declaration 'Old'.  Figure out how to resolve this situation,
@@ -3286,6 +3302,15 @@ bool Sema::MergeFunctionDecl(FunctionDecl *New, NamedDecl *&OldD,
     }
   }
 
+  if (New->hasAttr<ReqdWorkGroupSizeAttr>() &&
+      Old->hasAttr<ReqdWorkGroupSizeAttr>())
+    checkDimensionsAndSetDiagnostics<ReqdWorkGroupSizeAttr>(*this, New, Old);
+
+  if (New->hasAttr<SYCLIntelMaxWorkGroupSizeAttr>() &&
+      Old->hasAttr<SYCLIntelMaxWorkGroupSizeAttr>())
+    checkDimensionsAndSetDiagnostics<SYCLIntelMaxWorkGroupSizeAttr>(*this, New,
+                                                                    Old);
+
   if (New->hasAttr<InternalLinkageAttr>() &&
       !Old->hasAttr<InternalLinkageAttr>()) {
     Diag(New->getLocation(), diag::err_internal_linkage_redeclaration)

clang/lib/Sema/SemaSYCL.cpp

@@ -554,7 +554,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
 
       for (const CallGraphNode *CI : *N) {
         if (auto *Callee = dyn_cast<FunctionDecl>(CI->getDecl())) {
-          Callee = Callee->getCanonicalDecl();
+          Callee = Callee->getMostRecentDecl();
           if (!Visited.count(Callee))
             WorkList.push_back({Callee, FD});
         }

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

@@ -0,0 +1,68 @@
+// RUN: %clang_cc1 %s -I %S/Inputs -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -I %S/Inputs -fsyntax-only -fsycl -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify
+// RUN: %clang_cc1 %s -I %S/Inputs -fsyntax-only -ast-dump -fsycl -fsycl-is-device -triple spir64 -Wno-sycl-2017-compat | FileCheck %s
+
+#include <sycl.hpp>
+
+#ifndef TRIGGER_ERROR
+//first case - good case
+[[intelfpga::no_global_work_offset]] // expected-no-diagnostics
+void
+func1();
+
+[[intelfpga::max_work_group_size(4, 4, 4)]] void func1();
+
+[[cl::reqd_work_group_size(2, 2, 2)]] void func1() {}
+
+#else
+//second case - expect error
+[[intelfpga::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
+void
+func2();
+
+[[cl::reqd_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}}
+void
+func2() {}
+
+//third case - expect error
+[[cl::reqd_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
+void
+func3();
+
+[[cl::reqd_work_group_size(1, 1, 1)]] // expected-note {{conflicting attribute is here}}
+void
+// expected-warning@+1 {{attribute 'reqd_work_group_size' is already applied with different parameters}}
+func3() {} // expected-error {{'reqd_work_group_size' attribute conflicts with ''reqd_work_group_size'' attribute}}
+
+//fourth case - expect error
+[[intelfpga::max_work_group_size(4, 4, 4)]] // expected-note {{conflicting attribute is here}}
+void
+func4();
+
+[[intelfpga::max_work_group_size(8, 8, 8)]] // expected-note {{conflicting attribute is here}}
+void
+// expected-warning@+1 {{attribute 'max_work_group_size' is already applied with different parameters}}
+func4() {} // expected-error {{'max_work_group_size' attribute conflicts with ''max_work_group_size'' attribute}}
+#endif
+
+int main() {
+#ifndef TRIGGER_ERROR
+  // CHECK-LABEL:  FunctionDecl {{.*}} main 'int ()'
+  // CHECK:  `-FunctionDecl {{.*}}test_kernel1 'void ()'
+  // CHECK:  -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited 4 4 4
+  // CHECK:  -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} Inherited Enabled
+  // CHECK:  `-ReqdWorkGroupSizeAttr {{.*}} 2 2 2
+  cl::sycl::kernel_single_task<class test_kernel1>(
+      []() { func1(); });
+
+#else
+  cl::sycl::kernel_single_task<class test_kernel2>(
+      []() { func2(); }); // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
+
+  cl::sycl::kernel_single_task<class test_kernel3>(
+      []() { func3(); });
+
+  cl::sycl::kernel_single_task<class test_kernel4>(
+      []() { func4(); });
+#endif
+}