[SYCL] Provide a mechanism to silence incorrect sub group size warning (#11991)

As a follow up to #11687, this PR adds
a mechanism to silence the warning using dedicated switch
`-Wno-incorrect-sub-group-size` that is wrapped in the `-Wno-attribute`
group.

Author: Jakub Chlanda

clang/include/clang/Basic/DiagnosticGroups.td

@@ -780,9 +780,11 @@ def SyclIvdepAttribute : DiagGroup<"ivdep-compat">;
 def IndependentClassAttribute : DiagGroup<"IndependentClass-attribute">;
 def UnknownAttributes : DiagGroup<"unknown-attributes">;
 def IgnoredAttributes : DiagGroup<"ignored-attributes", [SyclIvdepAttribute]>;
+def IncorrectSubGroupSize: DiagGroup<"incorrect-sub-group-size">;
 def AcceptedAttributes : DiagGroup<"accepted-attributes">;
 def Attributes : DiagGroup<"attributes", [UnknownAttributes,
                                           IgnoredAttributes,
+                                          IncorrectSubGroupSize,
                                           AcceptedAttributes]>;
 def UnknownSanitizers : DiagGroup<"unknown-sanitizers">;
 def UnnamedTypeTemplateArgs : DiagGroup<"unnamed-type-template-args",

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -3539,7 +3539,7 @@ def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed"
 def warn_reqd_sub_group_attribute_n
     : Warning<"attribute argument %0 is invalid and will be ignored; %1 "
               "requires sub_group size %2">,
-  InGroup<IgnoredAttributes>;
+  InGroup<IncorrectSubGroupSize>;
 def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with"
   " exception specification; attribute ignored">,
   InGroup<IgnoredAttributes>;

clang/test/SemaSYCL/reqd-sub-group-size-silence-warnings.cpp

@@ -0,0 +1,24 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple nvptx -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-incorrect-sub-group-size %s
+// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify -Wno-attributes %s
+//
+// Sub group size of 8 is incompatible with both CUDA and HIP, expect it to be
+// silenced. Check both the dedicated switch '-Wno-incorrect-sub-group-size' and
+// the catch all '-Wno-attributes'.
+#include "sycl.hpp"
+
+
+// expected-no-diagnostics
+int main() {
+
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {});
+  });
+
+  return 0;
+}