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

[Yaxun Liu via cfe-commits]

https://github.com/yxsamliu created https://github.com/llvm/llvm-project/pull/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 -Woffload-incompatible-redeclare.

It won't cause warning in system headers, even with -Woffload-incompatible-redeclare.

>From 30daa345d23e91b66d66847cf329e056218c466e Mon Sep 17 00:00:00 2001
From: "Yaxun (Sam) Liu" <yaxun.liu at amd.com>
Date: Mon, 8 Jan 2024 14:02:06 -0500
Subject: [PATCH] [CUDA][HIP] warn incompatible redeclare

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 -Woffload-incompatible-redeclare.

It won't cause warning in system headers, even with
-Woffload-incompatible-redeclare.
---
 .../clang/Basic/DiagnosticSemaKinds.td        |  6 ++++
 clang/lib/Sema/SemaCUDA.cpp                   | 35 +++++++++++--------
 clang/test/SemaCUDA/function-redclare.cu      | 19 ++++++++++
 3 files changed, 46 insertions(+), 14 deletions(-)
 create mode 100644 clang/test/SemaCUDA/function-redclare.cu

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index a97182cad5d513..a4c5a76b59cb1b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8837,6 +8837,12 @@ 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<
+  "incompatible host/device attribute with redeclaration: "
+  "new declaration is %select{__device__|__global__|__host__|__host__ __device__}0 function, "
+  "old declaration is %select{__device__|__global__|__host__|__host__ __device__}1 function. "
+  "It will cause warning with nvcc">,
+  InGroup<DiagGroup<"offload-incompatible-redeclare">>, 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 6a66ecf6f94c17..27da1775d4751d 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -992,22 +992,29 @@ void Sema::checkCUDATargetOverload(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 == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(NewFD) &&
-            OldTarget == CFT_Device)) ||
-         (OldTarget == CFT_HostDevice &&
-          !(LangOpts.OffloadImplicitHostDeviceTemplates &&
-            isCUDAImplicitHostDeviceFunction(OldFD) &&
-            NewTarget == CFT_Device)) ||
-         (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) &&
         !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
                     /* ConsiderCudaAttrs = */ false)) {
-      Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
-          << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
-      Diag(OldFD->getLocation(), diag::note_previous_declaration);
-      NewFD->setInvalidDecl();
-      break;
+      if ((NewTarget == CFT_HostDevice &&
+           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+             isCUDAImplicitHostDeviceFunction(NewFD) &&
+             OldTarget == CFT_Device)) ||
+          (OldTarget == CFT_HostDevice &&
+           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
+             isCUDAImplicitHostDeviceFunction(OldFD) &&
+             NewTarget == CFT_Device)) ||
+          (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) {
+        Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
+            << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
+        Diag(OldFD->getLocation(), diag::note_previous_declaration);
+        NewFD->setInvalidDecl();
+        break;
+      }
+      if ((NewTarget == CFT_Host && OldTarget == CFT_Device) ||
+          (NewTarget == CFT_Device && OldTarget == CFT_Host)) {
+        Diag(NewFD->getLocation(), diag::warn_offload_incompatible_redeclare)
+            << NewTarget << 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 00000000000000..266483ebf49f0e
--- /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 -Woffload-incompatible-redeclare %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:   -isystem %S/Inputs -fcuda-is-device -Woffload-incompatible-redeclare -verify=redecl %s
+
+// expected-no-diagnostics
+#include "cuda.h"
+
+__device__ void f(); // redecl-note {{previous declaration is here}}
+
+void f() {} // redecl-warning {{incompatible host/device attribute with redeclaration: new declaration is __host__ function, old declaration is __device__ function. It will cause warning with nvcc}}
+
+void g(); // redecl-note {{previous declaration is here}}
+
+__device__ void g() {} // redecl-warning {{incompatible host/device attribute with redeclaration: new declaration is __device__ function, old declaration is __host__ function. It will cause warning with nvcc}}