[clang] [CUDA/HIP] Fix errors for device function used in host global initializers (PR #222338)
Mariya Podchishchaeva via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 10 07:46:52 PDT 2026
https://github.com/Fznamznon updated https://github.com/llvm/llvm-project/pull/222338
>From 87c539b9a77bfdd0e2401a710b95bd8c985d5efc Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <Mariya.Podchishchaeva at amd.com>
Date: Wed, 9 Sep 2026 09:14:13 -0500
Subject: [PATCH 1/4] [CUDA/HIP] Fix errors for device function used in host
global initializers
The existing checker was simply looking for a call or CXXConstructExpr which
would miss a call to device function in case an implicit cast was in place
or if a call was a part of an expression. Use a visitor to improve the
situation.
---
clang/lib/Sema/SemaCUDA.cpp | 54 +++++++++++++---------
clang/test/SemaCUDA/global-initializers.cu | 47 ++++++++++++++++++-
2 files changed, 78 insertions(+), 23 deletions(-)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13aa06aa97399..13d98b344cd97 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -776,39 +776,51 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
} else {
// This is a host-side global variable. Check that the initializer is
// callable from the host side.
- const FunctionDecl *InitFn = nullptr;
- if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
- InitFn = CE->getConstructor();
- } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
- InitFn = CE->getDirectCallee();
- }
- if (InitFn) {
- CUDAFunctionTarget InitFnTarget = IdentifyTarget(InitFn);
- if (InitFnTarget != CUDAFunctionTarget::Host &&
- InitFnTarget != CUDAFunctionTarget::HostDevice) {
- Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
- << InitFnTarget << InitFn;
- Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
- VD->setInvalidDecl();
- }
- }
+
struct GlobVarInitChecker : ConstEvaluatedExprVisitor<GlobVarInitChecker> {
+ private:
using Base = ConstEvaluatedExprVisitor<GlobVarInitChecker>;
SemaCUDA &SCRef;
- SourceLocation InitLoc;
+ VarDecl *VD;
+ void CheckForWrongSidedCall(const FunctionDecl *FD) {
+ CUDAFunctionTarget InitFnTarget = SCRef.IdentifyTarget(FD);
+ if (InitFnTarget != CUDAFunctionTarget::Host &&
+ InitFnTarget != CUDAFunctionTarget::HostDevice) {
+ SCRef.Diag(VD->getLocation(),
+ diag::err_ref_bad_target_global_initializer)
+ << InitFnTarget << FD;
+ SCRef.Diag(FD->getLocation(), diag::note_previous_decl) << FD;
+ VD->setInvalidDecl();
+ }
+ }
- GlobVarInitChecker(SemaCUDA &S, SourceLocation L)
- : Base(S.getASTContext()), SCRef(S), InitLoc(L) {}
+ public:
+ GlobVarInitChecker(SemaCUDA &S, VarDecl *VD)
+ : Base(S.getASTContext()), SCRef(S), VD(VD) {}
void VisitDeclRefExpr(const DeclRefExpr *DRE) {
if (auto *VarD = dyn_cast<VarDecl>(DRE->getDecl());
VarD && VarD->hasAttr<HIPManagedAttr>()) {
SCRef.Diag(DRE->getLocation(),
diag::err_cuda_invalid_use_of_managedvar);
- SCRef.Diag(InitLoc, diag::note_cuda_managed_var_in_glob_init);
+ SCRef.Diag(VD->getLocation(),
+ diag::note_cuda_managed_var_in_glob_init);
}
+ Base::VisitDeclRefExpr(DRE);
+ }
+ void VisitCallExpr(const CallExpr *CE) {
+ const FunctionDecl *InitFn = CE->getDirectCallee();
+ if (InitFn)
+ CheckForWrongSidedCall(InitFn);
+ Base::VisitCallExpr(CE);
+ }
+ void VisitCXXConstructExpr(const CXXConstructExpr *CE) {
+ const FunctionDecl *InitFn = CE->getConstructor();
+ if (InitFn)
+ CheckForWrongSidedCall(InitFn);
+ Base::VisitCXXConstructExpr(CE);
}
};
- GlobVarInitChecker Checker(*this, VD->getLocation());
+ GlobVarInitChecker Checker(*this, VD);
Checker.Visit(Init);
}
}
diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu
index 29e386134a3dd..3cd46b90638c7 100644
--- a/clang/test/SemaCUDA/global-initializers.cu
+++ b/clang/test/SemaCUDA/global-initializers.cu
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify
-// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify
+// RUN: %clang_cc1 %s -triple x86_64-linux-unknown -fsyntax-only -o - -verify=expected,host
+// RUN: %clang_cc1 %s -fcuda-is-device -triple nvptx -fsyntax-only -o - -verify=expected,device
#include "Inputs/cuda.h"
@@ -70,3 +70,46 @@ __device__ double AY = a.pow(2.0, 2); // expected-error{{dynamic initialization
const A ca;
const double CAX = ca.cpow(1.0, 1);
const __device__ double CAY = ca.cpow(2.0, 2);
+
+namespace ns1 {
+ // host-note at +3 {{'value_func' declared here}}
+ // expected-note at +2 5{{'value_func' declared here}}
+ // expected-note at +1 {{candidate function not viable: call to __device__ function from __host__ function}}
+__device__ constexpr inline int value_func() {
+ return 32;
+}
+}
+
+namespace ns2 {
+ using namespace ns1;
+ // diagnosed via overloading.
+ constexpr static unsigned var0 = value_func();
+ // expected-error at -1 {{no matching function for call to 'value_func'}}
+
+ // diagnosed via SemaCUDA::checkAllowedInitializer
+ constexpr static unsigned var1 = ns1::value_func();
+ // host-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ // FIXME: Inconsistency with var1 - non constexpr cases are diagnosed for both host and device.
+ static int var2 = 1 + ns1::value_func();
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+ static int var3 {ns1::value_func()};
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ struct A {
+ unsigned b;
+ };
+ A b{ns1::value_func()};
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ int foo(int);
+ int nested = foo(ns1::value_func());
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ void foobar() {
+ // diagnosed via SemaCUDA::checkCall
+ static int var2 = 1 + ns1::value_func();
+ // host-error at -1 {{reference to __device__ function 'value_func' in __host__ function}}
+ // device-error at -2 {{reference to __device__ function 'value_func' in global initializer}}
+ }
+}
>From 951d61aeeb57b4a9beac836b7e9aa22c445a7e9c Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <Mariya.Podchishchaeva at amd.com>
Date: Thu, 10 Sep 2026 07:46:52 -0500
Subject: [PATCH 2/4] Remove noop call
---
clang/lib/Sema/SemaCUDA.cpp | 1 -
1 file changed, 1 deletion(-)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13d98b344cd97..942f51fb52d5b 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -805,7 +805,6 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
SCRef.Diag(VD->getLocation(),
diag::note_cuda_managed_var_in_glob_init);
}
- Base::VisitDeclRefExpr(DRE);
}
void VisitCallExpr(const CallExpr *CE) {
const FunctionDecl *InitFn = CE->getDirectCallee();
>From 16aecda6948fd41034a3fb73e6a1e50033054d24 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <Mariya.Podchishchaeva at amd.com>
Date: Thu, 10 Sep 2026 09:21:30 -0500
Subject: [PATCH 3/4] Handle default initializers and default arguments
---
clang/lib/Sema/SemaCUDA.cpp | 19 ++++++++++++++++---
clang/test/SemaCUDA/global-initializers.cu | 16 +++++++++++++++-
2 files changed, 31 insertions(+), 4 deletions(-)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 942f51fb52d5b..e00b274adc2f4 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -806,18 +806,31 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
diag::note_cuda_managed_var_in_glob_init);
}
}
+
void VisitCallExpr(const CallExpr *CE) {
const FunctionDecl *InitFn = CE->getDirectCallee();
if (InitFn)
CheckForWrongSidedCall(InitFn);
Base::VisitCallExpr(CE);
}
+
void VisitCXXConstructExpr(const CXXConstructExpr *CE) {
- const FunctionDecl *InitFn = CE->getConstructor();
- if (InitFn)
- CheckForWrongSidedCall(InitFn);
+ const CXXConstructorDecl *Ctor = CE->getConstructor();
+ if (Ctor) {
+ CheckForWrongSidedCall(Ctor);
+ for (auto *I : Ctor->inits())
+ Visit(I->getInit());
+ }
Base::VisitCXXConstructExpr(CE);
}
+
+ void VisitCXXDefaultArgExpr(const CXXDefaultArgExpr *E) {
+ Visit(E->getExpr());
+ }
+
+ void VisitCXXDefaultInitExpr(const CXXDefaultInitExpr *E) {
+ Visit(E->getExpr());
+ }
};
GlobVarInitChecker Checker(*this, VD);
Checker.Visit(Init);
diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu
index 3cd46b90638c7..ef7d16f608ab0 100644
--- a/clang/test/SemaCUDA/global-initializers.cu
+++ b/clang/test/SemaCUDA/global-initializers.cu
@@ -73,7 +73,7 @@ const __device__ double CAY = ca.cpow(2.0, 2);
namespace ns1 {
// host-note at +3 {{'value_func' declared here}}
- // expected-note at +2 5{{'value_func' declared here}}
+ // expected-note at +2 7{{'value_func' declared here}}
// expected-note at +1 {{candidate function not viable: call to __device__ function from __host__ function}}
__device__ constexpr inline int value_func() {
return 32;
@@ -112,4 +112,18 @@ namespace ns2 {
// host-error at -1 {{reference to __device__ function 'value_func' in __host__ function}}
// device-error at -2 {{reference to __device__ function 'value_func' in global initializer}}
}
+
+ struct DefInit {
+ unsigned a = 1 + value_func();
+ };
+ DefInit testDefInit;
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ struct DefArg {
+ int data;
+ DefArg(int a = 1 + value_func()) : data(a) {}
+ };
+ DefArg testDefArg;
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
}
>From 1f7378fb2b23b84a5c6e88ad780d2ef2f973a3db Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <Mariya.Podchishchaeva at amd.com>
Date: Thu, 10 Sep 2026 09:39:59 -0500
Subject: [PATCH 4/4] Diagnose called function only once
---
clang/lib/Sema/SemaCUDA.cpp | 4 ++--
clang/test/SemaCUDA/global-initializers.cu | 11 ++++++++++-
2 files changed, 12 insertions(+), 3 deletions(-)
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index e00b274adc2f4..090030ea82503 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -785,7 +785,8 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
void CheckForWrongSidedCall(const FunctionDecl *FD) {
CUDAFunctionTarget InitFnTarget = SCRef.IdentifyTarget(FD);
if (InitFnTarget != CUDAFunctionTarget::Host &&
- InitFnTarget != CUDAFunctionTarget::HostDevice) {
+ InitFnTarget != CUDAFunctionTarget::HostDevice &&
+ !VD->isInvalidDecl()) {
SCRef.Diag(VD->getLocation(),
diag::err_ref_bad_target_global_initializer)
<< InitFnTarget << FD;
@@ -806,7 +807,6 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
diag::note_cuda_managed_var_in_glob_init);
}
}
-
void VisitCallExpr(const CallExpr *CE) {
const FunctionDecl *InitFn = CE->getDirectCallee();
if (InitFn)
diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu
index ef7d16f608ab0..cba5c539614c7 100644
--- a/clang/test/SemaCUDA/global-initializers.cu
+++ b/clang/test/SemaCUDA/global-initializers.cu
@@ -73,11 +73,14 @@ const __device__ double CAY = ca.cpow(2.0, 2);
namespace ns1 {
// host-note at +3 {{'value_func' declared here}}
- // expected-note at +2 7{{'value_func' declared here}}
+ // expected-note at +2 9{{'value_func' declared here}}
// expected-note at +1 {{candidate function not viable: call to __device__ function from __host__ function}}
__device__ constexpr inline int value_func() {
return 32;
}
+__device__ constexpr inline int another_value_func() {
+ return 32;
+}
}
namespace ns2 {
@@ -126,4 +129,10 @@ namespace ns2 {
DefArg testDefArg;
// expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+ unsigned twotimes = ns1::value_func() + ns1::another_value_func();
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
+ unsigned twotimes1 = ns1::value_func() + ns1::value_func();
+ // expected-error at -1 {{reference to __device__ function 'value_func' in global initializer}}
+
}
More information about the cfe-commits
mailing list