Early return and merge warning

Author: jchlanda

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -3439,10 +3439,6 @@ def err_attribute_argument_is_zero : Error<
 def warn_attribute_argument_n_negative : Warning<
   "%0 attribute parameter %1 is negative and will be ignored">,
   InGroup<CudaCompat>;
-def warn_reqd_sub_group_attribute_cuda_n_32
-    : Warning<"attribute argument %0 is invalid and will be ignored; CUDA "
-              "requires sub_group size 32">,
-      InGroup<CudaCompat>;
 def err_property_function_in_objc_container : Error<
   "use of Objective-C property in function nested in Objective-C "
   "container not supported, move function outside its container">;
@@ -3541,9 +3537,9 @@ def warn_dllimport_dropped_from_inline_function : Warning<
 def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed"
   " only on a function directly called from a SYCL kernel function; attribute ignored">,
   InGroup<IgnoredAttributes>;
-def warn_amd_reqd_sub_group_attribute_n
-    : Warning<"attribute argument %0 is invalid and will be ignored; AMD "
-              "requires sub_group size %1">,
+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>;
 def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with"
   " exception specification; attribute ignored">,

clang/lib/Sema/SemaDeclAttr.cpp

@@ -4009,8 +4009,9 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
     }
     auto &TI = Context.getTargetInfo();
     if (TI.getTriple().isNVPTX() && ArgVal != 32) {
-      Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_cuda_n_32)
-          << ArgVal.getSExtValue();
+      Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n)
+          << ArgVal.getSExtValue() << TI.getTriple().getArchName() << 32;
+      return;
     }
     if (TI.getTriple().isAMDGPU()) {
       const auto HasWaveFrontSize64 =
@@ -4024,8 +4025,10 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI,
       const auto SupportedWaveFrontSize =
           HasWaveFrontSize64 && !HasWaveFrontSize32 ? 64 : 32;
       if (ArgVal != SupportedWaveFrontSize) {
-        Diag(E->getExprLoc(), diag::warn_amd_reqd_sub_group_attribute_n)
-            << ArgVal.getSExtValue() << SupportedWaveFrontSize;
+        Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n)
+            << ArgVal.getSExtValue() << TI.getTriple().getArchName()
+            << SupportedWaveFrontSize;
+        return;
       }
     }
 

clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp

@@ -9,15 +9,15 @@ int main() {
   sycl::queue Q;
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; AMD requires sub_group size 32}}
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; amdgcn requires sub_group size 32}}
   });
 
   Q.submit([&](sycl::handler &h) {
     h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(32)]] {});
   });
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 32}}
+    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 32}}
   });
 
   return 0;

clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp

@@ -13,11 +13,11 @@ int main() {
   });
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; AMD requires sub_group size 64}}
+    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; amdgcn requires sub_group size 64}}
   });
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 64}}
+    h.single_task<class valid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 64}}
   });
 
   return 0;

clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp

@@ -9,7 +9,7 @@ int main() {
   sycl::queue Q;
 
   Q.submit([&](sycl::handler &h) {
-    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; CUDA requires sub_group size 32}}
+    h.single_task<class invalid_kernel>([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; nvptx requires sub_group size 32}}
   });
 
   Q.submit([&](sycl::handler &h) {