[clang] [CUDA/HIP] Do not check function calls in discarded statement (PR #194606)

[Weibo He via cfe-commits]

https://github.com/NewSigma created https://github.com/llvm/llvm-project/pull/194606

Previously, calling a host-device mismatch function inside a discarded `if constexpr` branch would trigger an error. This patch recognizes that discarded statements are never instantiated and allows such code.

>From 576efbbf456f88731b3c5d36bfa4a70b338404fd Mon Sep 17 00:00:00 2001
From: NewSigma <NewSigma at 163.com>
Date: Sun, 26 Apr 2026 19:38:58 +0800
Subject: [PATCH] [SemaCUDA] Do not check function calls in discarded statement

---
 clang/lib/Sema/SemaCUDA.cpp                   |  3 ++-
 .../test/SemaCUDA/call-device-fn-from-host.cu | 19 +++++++++++++++++--
 .../test/SemaCUDA/call-host-fn-from-device.cu | 17 ++++++++++++++++-
 3 files changed, 35 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index c086f9a32ce4e..56db70d05761e 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -990,7 +990,8 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
   assert(Callee && "Callee may not be null.");
 
   const auto &ExprEvalCtx = SemaRef.currentEvaluationContext();
-  if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
+  if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated() ||
+      ExprEvalCtx.isDiscardedStatementContext())
     return true;
 
   // C++ deduction guides participate in overload resolution but are not
diff --git a/clang/test/SemaCUDA/call-device-fn-from-host.cu b/clang/test/SemaCUDA/call-device-fn-from-host.cu
index 4d66fccd84d53..7d72807494fa9 100644
--- a/clang/test/SemaCUDA/call-device-fn-from-host.cu
+++ b/clang/test/SemaCUDA/call-device-fn-from-host.cu
@@ -1,6 +1,6 @@
-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN: %clang_cc1 %s --std=c++17 -triple x86_64-unknown-linux -emit-llvm -o - \
 // RUN:   -verify -verify-ignore-unexpected=note
-// RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \
+// RUN: %clang_cc1 %s --std=c++17 -triple x86_64-unknown-linux -emit-llvm -o - \
 // RUN:   -verify=expected,omp -verify-ignore-unexpected=note -fopenmp
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
@@ -97,3 +97,18 @@ void host_func(void) { kernel<<<1, 1>>>(); }
 __device__ void f();
 template<void(*F)()> __global__ void t() { F(); }
 __host__ void g() { t<f><<<1,1>>>(); }
+
+namespace template_if_constexpr {
+  template<bool B>
+  __host__ __device__ void fn() {
+    if constexpr (!B)
+      device_fn();
+
+    if constexpr (B)
+      device_fn(); // expected-error {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
+  }
+
+  void call() {
+    fn<true>();
+  }
+}
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index acdd291b66457..6bcb6f2ff0bc4 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
+// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \
 // RUN:   -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note
 
 // Note: This test won't work with -fsyntax-only, because some of these errors
@@ -138,3 +138,18 @@ __host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
 // expected-error at -1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
 
 __device__ void double_specialization() { TmplStruct<int>().fn<int>(); }
+
+namespace template_if_constexpr {
+  template<bool B>
+  __host__ __device__ void fn() {
+    if constexpr (!B)
+      host_fn();
+
+    if constexpr (B)
+      host_fn(); // expected-error {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
+  }
+
+  __device__ void call() {
+    fn<true>();
+  }
+}