[clang] [CUDA/HIP] Do not check function calls in discarded statement (PR #194606)
Weibo He via cfe-commits
cfe-commits at lists.llvm.org
Mon May 4 07:11:26 PDT 2026
https://github.com/NewSigma updated https://github.com/llvm/llvm-project/pull/194606
>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 1/5] [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>();
+ }
+}
>From 21305aea7353020703d9e327e66ee2d6fbfd39f5 Mon Sep 17 00:00:00 2001
From: NewSigma <NewSigma at 163.com>
Date: Fri, 1 May 2026 17:51:18 +0800
Subject: [PATCH 2/5] Skip diagnostic for uninstantiated HD to global call
---
clang/lib/Sema/SemaCUDA.cpp | 17 ++++++++++++-----
clang/test/SemaCUDA/device-kernel-call.cu | 16 ++++++++++++++++
2 files changed, 28 insertions(+), 5 deletions(-)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 56db70d05761e..9e05de941f335 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -84,10 +84,6 @@ ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
return ExprError(
Diag(LLLLoc, diag::err_cuda_device_kernel_launch_not_supported));
- if (IsDeviceKernelCall && !getLangOpts().GPURelocatableDeviceCode)
- return ExprError(
- Diag(LLLLoc, diag::err_cuda_device_kernel_launch_require_rdc));
-
FunctionDecl *ConfigDecl = IsDeviceKernelCall
? getASTContext().getcudaLaunchDeviceDecl()
: getASTContext().getcudaConfigureCallDecl();
@@ -1027,9 +1023,20 @@ bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
}
}();
+ bool IsDeviceKernelCall = Callee == getASTContext().getcudaLaunchDeviceDecl();
+ bool CallerHD = Caller && Caller->hasAttr<CUDAHostAttr>() &&
+ Caller->hasAttr<CUDADeviceAttr>();
+ bool CallerDiscard = SemaRef.getEmissionStatus(Caller) ==
+ Sema::FunctionEmissionStatus::TemplateDiscarded;
+ bool RDC = getLangOpts().GPURelocatableDeviceCode;
+ if (IsDeviceKernelCall && !(CallerHD && CallerDiscard) && !RDC) {
+ Diag(Loc, diag::err_cuda_device_kernel_launch_require_rdc);
+ return false;
+ }
+
if (DiagKind == SemaDiagnosticBuilder::K_Nop) {
// For -fgpu-rdc, keep track of external kernels used by host functions.
- if (getLangOpts().CUDAIsDevice && getLangOpts().GPURelocatableDeviceCode &&
+ if (getLangOpts().CUDAIsDevice && RDC &&
Callee->hasAttr<CUDAGlobalAttr>() && !Callee->isDefined() &&
(!Caller || (!Caller->getDescribedFunctionTemplate() &&
getASTContext().GetGVALinkageForFunction(Caller) ==
diff --git a/clang/test/SemaCUDA/device-kernel-call.cu b/clang/test/SemaCUDA/device-kernel-call.cu
index 856cbd88404e6..7511cf148a077 100644
--- a/clang/test/SemaCUDA/device-kernel-call.cu
+++ b/clang/test/SemaCUDA/device-kernel-call.cu
@@ -13,3 +13,19 @@ __global__ void g1(void) {
// nordc-error at -1 {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}
// hip-error at -2 {{device-side kernel call/launch is not supported}}
}
+
+namespace template_if_constexpr {
+ template<bool B>
+ __host__ __device__ void fn() {
+ if constexpr (B)
+ g2<<<1, 1>>>(42);
+ // hip-error at -1 {{device-side kernel call/launch is not supported}}
+ }
+
+ void call() {
+ fn<false>();
+ fn<true>();
+ // nordc-error at -7 {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}
+ // nordc-note at -2 {{in instantiation of function template specialization 'template_if_constexpr::fn<true>' requested here}}
+ }
+}
>From ebbb460eebbfdeb358d5c32ae739bfb234152c6a Mon Sep 17 00:00:00 2001
From: NewSigma <NewSigma at 163.com>
Date: Fri, 1 May 2026 20:08:30 +0800
Subject: [PATCH 3/5] Update tests
---
clang/test/SemaCUDA/call-device-fn-from-host.cu | 8 +++-----
clang/test/SemaCUDA/call-host-fn-from-device.cu | 8 +++-----
2 files changed, 6 insertions(+), 10 deletions(-)
diff --git a/clang/test/SemaCUDA/call-device-fn-from-host.cu b/clang/test/SemaCUDA/call-device-fn-from-host.cu
index 7d72807494fa9..cb8f29e3d9499 100644
--- a/clang/test/SemaCUDA/call-device-fn-from-host.cu
+++ b/clang/test/SemaCUDA/call-device-fn-from-host.cu
@@ -101,14 +101,12 @@ __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}}
+ device_fn();
}
void call() {
- fn<true>();
+ fn<false>();
+ fn<true>(); // expected-error at -5 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
}
}
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index 6bcb6f2ff0bc4..2b42f4c4d75fc 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -142,14 +142,12 @@ __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}}
+ host_fn();
}
__device__ void call() {
- fn<true>();
+ fn<false>();
+ fn<true>(); // expected-error at -5 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
}
}
>From fef4399a7220fbc0064c522dee4fcc48f68ed609 Mon Sep 17 00:00:00 2001
From: NewSigma <NewSigma at 163.com>
Date: Mon, 4 May 2026 18:44:12 +0800
Subject: [PATCH 4/5] Add release note
---
clang/docs/ReleaseNotes.rst | 3 +++
1 file changed, 3 insertions(+)
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index ec64c2008d89b..37514a55e56f3 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -470,6 +470,9 @@ Improvements to Clang's diagnostics
- Removed the body of lambdas from some diagnostic messages.
+- Fixed false positive host-device mismatch errors in discarded `if constexpr` branches for CUDA/HIP;
+ such calls are now correctly skipped.
+
Improvements to Clang's time-trace
----------------------------------
>From 419bf5ac93ad07745193fd72130b8ff555885347 Mon Sep 17 00:00:00 2001
From: NewSigma <NewSigma at 163.com>
Date: Mon, 4 May 2026 21:01:59 +0800
Subject: [PATCH 5/5] Keep --std=c++11 checks
---
clang/test/SemaCUDA/call-device-fn-from-host.cu | 7 +++++++
clang/test/SemaCUDA/call-host-fn-from-device.cu | 5 +++++
2 files changed, 12 insertions(+)
diff --git a/clang/test/SemaCUDA/call-device-fn-from-host.cu b/clang/test/SemaCUDA/call-device-fn-from-host.cu
index cb8f29e3d9499..64394c7a4d958 100644
--- a/clang/test/SemaCUDA/call-device-fn-from-host.cu
+++ b/clang/test/SemaCUDA/call-device-fn-from-host.cu
@@ -1,3 +1,8 @@
+// RUN: %clang_cc1 %s --std=c++11 -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: -verify=expected,omp -verify-ignore-unexpected=note -fopenmp
+
// 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++17 -triple x86_64-unknown-linux -emit-llvm -o - \
@@ -98,6 +103,7 @@ __device__ void f();
template<void(*F)()> __global__ void t() { F(); }
__host__ void g() { t<f><<<1,1>>>(); }
+#if __cplusplus >= 201703L
namespace template_if_constexpr {
template<bool B>
__host__ __device__ void fn() {
@@ -110,3 +116,4 @@ namespace template_if_constexpr {
fn<true>(); // expected-error at -5 {{reference to __device__ function 'device_fn' in __host__ __device__ function}}
}
}
+#endif
diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu
index 2b42f4c4d75fc..d172cd966c823 100644
--- a/clang/test/SemaCUDA/call-host-fn-from-device.cu
+++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu
@@ -1,3 +1,6 @@
+// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \
+// RUN: -emit-llvm -o /dev/null -verify -verify-ignore-unexpected=note
+
// 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
@@ -139,6 +142,7 @@ __host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }
+#if __cplusplus >= 201703L
namespace template_if_constexpr {
template<bool B>
__host__ __device__ void fn() {
@@ -151,3 +155,4 @@ namespace template_if_constexpr {
fn<true>(); // expected-error at -5 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
}
}
+#endif
More information about the cfe-commits
mailing list