[clang] 53d2f4d - [CUDA][HIP] warn incompatible redeclare (#77359)

[via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2024-06-10T10:08:26-04:00
New Revision: 53d2f4d967838468423715178a8344739d7a63c9

URL: https://github.com/llvm/llvm-project/commit/53d2f4d967838468423715178a8344739d7a63c9
DIFF: https://github.com/llvm/llvm-project/commit/53d2f4d967838468423715178a8344739d7a63c9.diff

LOG: [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.

Added: 
    clang/test/SemaCUDA/function-redclare.cu

Modified: 
    clang/include/clang/Basic/DiagnosticSemaKinds.td
    clang/lib/Sema/SemaCUDA.cpp
    llvm/docs/CompileCudaWithLLVM.rst

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index de3aa4b4028eb..193eae3bc41d6 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9015,6 +9015,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 "

diff  --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 80ea43dc5316e..580b9872c6a1d 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);
+      }
     }
   }
 }

diff  --git a/clang/test/SemaCUDA/function-redclare.cu b/clang/test/SemaCUDA/function-redclare.cu
new file mode 100644
index 0000000000000..7cd9bad79ae98
--- /dev/null
+++ b/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}}

diff  --git a/llvm/docs/CompileCudaWithLLVM.rst b/llvm/docs/CompileCudaWithLLVM.rst
index 631691ef9b472..0371d7a3bdfcb 100644
--- a/llvm/docs/CompileCudaWithLLVM.rst
+++ b/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
 --------------------------------------