Update PropagateAndDiagnoseDeviceAttr

intel · Feb 28, 2022 · 3c5c444 · 3c5c444
1 parent 5410e23
commit 3c5c444
Show file tree

Hide file tree

Showing 2 changed files with 35 additions and 1 deletion.
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
@@ -3927,6 +3927,22 @@ static void PropagateAndDiagnoseDeviceAttr(
     }
     break;
   }
+  case attr::Kind::WorkGroupSizeHint: {
+    auto *WGSH = cast<WorkGroupSizeHintAttr>(A);
+    if (auto *Existing = SYCLKernel->getAttr<WorkGroupSizeHintAttr>()) {
+      if (Existing->getXDimVal() != WGSH->getXDimVal() ||
+          Existing->getYDimVal() != WGSH->getYDimVal() ||
+          Existing->getZDimVal() != WGSH->getZDimVal()) {
+        S.Diag(SYCLKernel->getLocation(),
+               diag::err_conflicting_sycl_kernel_attributes);
+        S.Diag(Existing->getLocation(), diag::note_conflicting_attribute);
+        S.Diag(WGSH->getLocation(), diag::note_conflicting_attribute);
+        SYCLKernel->setInvalidDecl();
+      }
+    }
+    SYCLKernel->addAttr(A);
+    break;
+  }
   case attr::Kind::SYCLIntelMaxWorkGroupSize: {
     auto *SIMWGSA = cast<SYCLIntelMaxWorkGroupSizeAttr>(A);
     if (auto *Existing = SYCLKernel->getAttr<ReqdWorkGroupSizeAttr>()) {

diff --git a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp
@@ -105,11 +105,25 @@ class Functor4x4x4 {
   [[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
 };
 
-[[sycl::work_group_size_hint(4, 4, 4)]] void f4x4x4(){};
+#if defined(EXPECT_PROP) && defined(TRIGGER_ERROR)
+[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; // expected-note {{conflicting attribute is here}}
+
+class FunctorConflict { // expected-error {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}}
+public:
+  [[sycl::work_group_size_hint(16, 2, 1)]] void operator()() const { // expected-note {{conflicting attribute is here}}
+    f8x8x8();
+  };
+};
+#endif
 
 void invoke() {
   Functor16x2x1 f16x2x1;
   Functor4x4x4 f4x4x4;
+
+#if defined(EXPECT_PROP) && defined(TRIGGER_ERROR)
+  FunctorConflict fConflict;
+#endif
+
   sycl::queue q;
 
   q.submit([&](sycl::handler &h) {
@@ -146,6 +160,10 @@ void invoke() {
     // CHECK-NEXT:  ConstantExpr{{.*}}'int'
     // CHECK-NEXT:  value: Int 4
     // CHECK-NEXT:  IntegerLiteral{{.*}}4{{$}}
+
+#if defined(EXPECT_PROP) && defined(TRIGGER_ERROR)
+    h.single_task<class kernel_3>(fConflict);
+#endif
   });
 
   // FIXME: Add tests with the C++23 lambda attribute syntax.