[clang] [Clang] Fix sema checks thinking kernels aren't kernels (PR #104460)

[Artem Belevich via cfe-commits]

================
@@ -7163,24 +7163,27 @@ void Sema::ProcessDeclAttributeList(
     } else if (const auto *A = D->getAttr<OpenCLIntelReqdSubGroupSizeAttr>()) {
       Diag(D->getLocation(), diag::err_opencl_kernel_attr) << A;
       D->setInvalidDecl();
-    } else if (!D->hasAttr<CUDAGlobalAttr>()) {
-      if (const auto *A = D->getAttr<AMDGPUFlatWorkGroupSizeAttr>()) {
-        Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
-            << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction;
-        D->setInvalidDecl();
-      } else if (const auto *A = D->getAttr<AMDGPUWavesPerEUAttr>()) {
-        Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
-            << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction;
-        D->setInvalidDecl();
-      } else if (const auto *A = D->getAttr<AMDGPUNumSGPRAttr>()) {
-        Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
-            << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction;
-        D->setInvalidDecl();
-      } else if (const auto *A = D->getAttr<AMDGPUNumVGPRAttr>()) {
-        Diag(D->getLocation(), diag::err_attribute_wrong_decl_type)
-            << A << A->isRegularKeywordAttribute() << ExpectedKernelFunction;
-        D->setInvalidDecl();
-      }
+    }
+  }
+  const FunctionType *FnTy = D->getFunctionType();
+  if (!D->hasAttr<CUDAGlobalAttr>() && !D->hasAttr<OpenCLKernelAttr>() &&
+      FnTy && FnTy->getCallConv() != CallingConv::CC_AMDGPUKernelCall) {
----------------
Artem-B wrote:

Nit: Way too many negations. My favorite pet peeve.

`if (!( __global__ || OpenCL || AMDGPUKernelCall))` would be easier to understand.

https://github.com/llvm/llvm-project/pull/104460