[CUDA][HIP] warn incompatible redeclare (#77359)

nvcc warns about the following code:

`void f();
__device__ void f() {}`

but clang does not since clang allows device function to overload host
function.

Users want clang to emit similar warning to help code to be compatible
with nvcc.

Since this may cause regression with existing code, the warning is off
by default and can be enabled by -Wnvcc-compat.

It won't cause warning in system headers, even with -Wnvcc-compat.

1 files changed, 25 insertions, 16 deletions

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 80ea43d..580b987 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -1018,24 +1018,33 @@ void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
     // HD/global functions "exist" in some sense on both the host and device, so
     // should have the same implementation on both sides.
     if (NewTarget != OldTarget &&
-        ((NewTarget == CUDAFunctionTarget::HostDevice &&
-          !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
-            isImplicitHostDeviceFunction(NewFD) &&
-            OldTarget == CUDAFunctionTarget::Device)) ||
-         (OldTarget == CUDAFunctionTarget::HostDevice &&
-          !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
-            isImplicitHostDeviceFunction(OldFD) &&
-            NewTarget == CUDAFunctionTarget::Device)) ||
-         (NewTarget == CUDAFunctionTarget::Global) ||
-         (OldTarget == CUDAFunctionTarget::Global)) &&
         !SemaRef.IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
                             /* ConsiderCudaAttrs = */ false)) {
-      Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
-          << llvm::to_underlying(NewTarget) << NewFD->getDeclName()
-          << llvm::to_underlying(OldTarget) << OldFD;
-      Diag(OldFD->getLocation(), diag::note_previous_declaration);
-      NewFD->setInvalidDecl();
-      break;
+      if ((NewTarget == CUDAFunctionTarget::HostDevice &&
+           !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
+             isImplicitHostDeviceFunction(NewFD) &&
+             OldTarget == CUDAFunctionTarget::Device)) ||
+          (OldTarget == CUDAFunctionTarget::HostDevice &&
+           !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
+             isImplicitHostDeviceFunction(OldFD) &&
+             NewTarget == CUDAFunctionTarget::Device)) ||
+          (NewTarget == CUDAFunctionTarget::Global) ||
+          (OldTarget == CUDAFunctionTarget::Global)) {
+        Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
+            << llvm::to_underlying(NewTarget) << NewFD->getDeclName()
+            << llvm::to_underlying(OldTarget) << OldFD;
+        Diag(OldFD->getLocation(), diag::note_previous_declaration);
+        NewFD->setInvalidDecl();
+        break;
+      }
+      if ((NewTarget == CUDAFunctionTarget::Host &&
+           OldTarget == CUDAFunctionTarget::Device) ||
+          (NewTarget == CUDAFunctionTarget::Device &&
+           OldTarget == CUDAFunctionTarget::Host)) {
+        Diag(NewFD->getLocation(), diag::warn_offload_incompatible_redeclare)
+            << llvm::to_underlying(NewTarget) << llvm::to_underlying(OldTarget);
+        Diag(OldFD->getLocation(), diag::note_previous_declaration);
+      }
     }
   }
 }