[clang] [HIP] Suppress unused warnings for wrong-side overload candidates (PR #195714)

[via cfe-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-clang

Author: Yaxun (Sam) Liu (yxsamliu)

<details>
<summary>Changes</summary>

A HIP source file is checked twice: once for the host and once for the device.
When an implicitly host/device template does overload resolution, each pass may
pick the overload that belongs to that side.

For example, a host/device template may call `operator==`. The host pass picks
the host-only `operator==`, while the device pass picks the device-only
`operator==`. The overload from the other side can then look unused in the
current pass, so `-Wunused-function -Werror` rejects code that is valid for the
other pass.

Track CUDA/HIP wrong-side overload candidates during overload resolution and
suppress the current-side unused-function warning for those declarations. This
keeps the warning for normal unused functions, but avoids false positives for
host-only helpers seen during device compilation and device-only helpers seen
during host compilation.



---
Full diff: https://github.com/llvm/llvm-project/pull/195714.diff


4 Files Affected:

- (modified) clang/include/clang/AST/ASTContext.h (+5) 
- (modified) clang/lib/Sema/SemaDecl.cpp (+11) 
- (modified) clang/lib/Sema/SemaOverload.cpp (+10-3) 
- (added) clang/test/SemaCUDA/unused-wrong-side-overload-candidate.cu (+25) 


``````````diff
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index c45d54fdd2e88..a0c155616b380 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -1395,6 +1395,11 @@ class ASTContext : public RefCountedBase<ASTContext> {
   /// in device compilation.
   llvm::DenseSet<const FunctionDecl *> CUDAImplicitHostDeviceFunUsedByDevice;
 
+  /// Keep track of CUDA/HIP wrong-side overload candidates seen in overload
+  /// resolution. Their usage is owned by the opposite compilation side, so the
+  /// current side should not warn that they are unused.
+  llvm::DenseSet<const FunctionDecl *> CUDAWrongSideOverloadCandidates;
+
   /// Map of SYCL kernels indexed by the unique type used to name the kernel.
   /// Entries are not serialized but are recreated on deserialization of a
   /// sycl_kernel_entry_point attributed function declaration.
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index eb5b6d65b4d58..907d8855b515e 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -1927,6 +1927,17 @@ bool Sema::ShouldWarnIfUnusedFileScopedDecl(const DeclaratorDecl *D) const {
     return false;
 
   if (const FunctionDecl *FD = dyn_cast<FunctionDecl>(D)) {
+    if (LangOpts.CUDA && Context.CUDAWrongSideOverloadCandidates.contains(FD)) {
+      bool IsHost =
+          FD->hasAttr<CUDAHostAttr>() || !FD->hasAttr<CUDADeviceAttr>();
+      bool IsDeviceOnly =
+          !FD->hasAttr<CUDAHostAttr>() && FD->hasAttr<CUDADeviceAttr>();
+      bool IsGlobal = FD->hasAttr<CUDAGlobalAttr>();
+      if ((LangOpts.CUDAIsDevice && IsHost) ||
+          (!LangOpts.CUDAIsDevice && (IsDeviceOnly || IsGlobal)))
+        return false;
+    }
+
     if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation)
       return false;
     // A non-out-of-line declaration of a member specialization was implicitly
diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp
index 96c4ce489fe04..c737ba17b0329 100644
--- a/clang/lib/Sema/SemaOverload.cpp
+++ b/clang/lib/Sema/SemaOverload.cpp
@@ -7452,6 +7452,9 @@ void Sema::AddOverloadCandidate(
   // (CUDA B.1): Check for invalid calls between targets.
   if (getLangOpts().CUDA) {
     const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+    if (Caller && Function &&
+        CUDA().IdentifyPreference(Caller, Function) == SemaCUDA::CFP_WrongSide)
+      getASTContext().CUDAWrongSideOverloadCandidates.insert(Function);
     // Skip the check for callers that are implicit members, because in this
     // case we may not yet know what the member's target is; the target is
     // inferred for the member automatically, based on the bases and fields of
@@ -8022,13 +8025,17 @@ void Sema::AddMethodCandidate(
   }
 
   // (CUDA B.1): Check for invalid calls between targets.
-  if (getLangOpts().CUDA)
-    if (!CUDA().IsAllowedCall(getCurFunctionDecl(/*AllowLambda=*/true),
-                              Method)) {
+  if (getLangOpts().CUDA) {
+    const FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
+    if (Caller && Method &&
+        CUDA().IdentifyPreference(Caller, Method) == SemaCUDA::CFP_WrongSide)
+      getASTContext().CUDAWrongSideOverloadCandidates.insert(Method);
+    if (!CUDA().IsAllowedCall(Caller, Method)) {
       Candidate.Viable = false;
       Candidate.FailureKind = ovl_fail_bad_target;
       return;
     }
+  }
 
   if (Method->getTrailingRequiresClause()) {
     ConstraintSatisfaction Satisfaction;
diff --git a/clang/test/SemaCUDA/unused-wrong-side-overload-candidate.cu b/clang/test/SemaCUDA/unused-wrong-side-overload-candidate.cu
new file mode 100644
index 0000000000000..3d3015ca222a5
--- /dev/null
+++ b/clang/test/SemaCUDA/unused-wrong-side-overload-candidate.cu
@@ -0,0 +1,25 @@
+// RUN: %clang_cc1 -std=c++20 -x hip -fcuda-is-device \
+// RUN:   -foffload-implicit-host-device-templates -Wall -Werror \
+// RUN:   -fsyntax-only %s
+// RUN: %clang_cc1 -std=c++20 -x hip \
+// RUN:   -foffload-implicit-host-device-templates -Wall -Werror \
+// RUN:   -fsyntax-only %s
+
+#include "Inputs/cuda.h"
+
+struct BothTy {};
+
+static bool operator==(BothTy, BothTy) { return true; }
+static __device__ bool operator==(BothTy, BothTy) { return true; }
+
+template <class T> bool compare(T LHS, T RHS) {
+  return LHS == RHS;
+}
+
+__host__ bool host_use() {
+  return compare(BothTy{}, BothTy{});
+}
+
+__device__ bool device_use() {
+  return compare(BothTy{}, BothTy{});
+}

``````````

</details>


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