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

This PR was already merged in trunk:

llvm#77359

However, it was reverted in amd-staging due
to ck issue ROCm/composable_kernel#1330

Now try relanding it in amd-staging after ck issue fixed
by

ROCm/composable_kernel#1342

Fixes: SWDEV-431838

Original commit message:

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.

Change-Id: Ia370700eb3eb1b1928d04ec59e2ec63506f85545

Author: yxsamliu

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9067,6 +9067,11 @@ def err_cuda_ovl_target : Error<
   "cannot overload %select{__device__|__global__|__host__|__host__ __device__}2 function %3">;
 def note_cuda_ovl_candidate_target_mismatch : Note<
     "candidate template ignored: target attributes do not match">;
+def warn_offload_incompatible_redeclare : Warning<
+  "target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:"
+  "new declaration is %select{__device__|__global__|__host__|__host__ __device__}0 function, "
+  "old declaration is %select{__device__|__global__|__host__|__host__ __device__}1 function">,
+  InGroup<DiagGroup<"nvcc-compat">>, DefaultIgnore;
 
 def err_cuda_device_builtin_surftex_cls_template : Error<
     "illegal device builtin %select{surface|texture}0 reference "

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);
+      }
     }
   }
 }

clang/test/SemaCUDA/function-redclare.cu

@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:   -isystem %S/Inputs -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:   -isystem %S/Inputs -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:   -isystem %S/Inputs -verify=redecl -Wnvcc-compat %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:   -isystem %S/Inputs -fcuda-is-device -Wnvcc-compat -verify=redecl %s
+
+// expected-no-diagnostics
+#include "cuda.h"
+
+__device__ void f(); // redecl-note {{previous declaration is here}}
+
+void f() {} // redecl-warning {{target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:new declaration is __host__ function, old declaration is __device__ function}}
+
+void g(); // redecl-note {{previous declaration is here}}
+
+__device__ void g() {} // redecl-warning {{target-attribute based function overloads are not supported by NVCC and will be treated as a function redeclaration:new declaration is __device__ function, old declaration is __host__ function}}

llvm/docs/CompileCudaWithLLVM.rst

@@ -418,6 +418,17 @@ the compiler chooses to inline ``host_only``.
 Member functions, including constructors, may be overloaded using H and D
 attributes.  However, destructors cannot be overloaded.
 
+Clang Warnings for Host and Device Function Declarations
+--------------------------------------------------------
+
+Clang can emit warnings when it detects that host (H) and device (D) functions are declared or defined with the same signature. These warnings are not enabled by default.
+
+To enable these warnings, use the following compiler flag:
+
+.. code-block:: console
+
+    -Wnvcc-compat
+
 Using a Different Class on Host/Device
 --------------------------------------