[clang] 4735d35 - [CUDA/HIP] Fix errors for device function used in host global initializers (#222338)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Sep 10 23:34:03 PDT 2026
Author: Mariya Podchishchaeva
Date: 2026-09-11T08:33:58+02:00
New Revision: 4735d351182f874d61690f6fc2d2e3b154914ad5
URL: https://github.com/llvm/llvm-project/commit/4735d351182f874d61690f6fc2d2e3b154914ad5
DIFF: https://github.com/llvm/llvm-project/commit/4735d351182f874d61690f6fc2d2e3b154914ad5.diff
LOG: [CUDA/HIP] Fix errors for device function used in host global initializers (#222338)
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.
Added:
Modified:
clang/lib/Sema/SemaCUDA.cpp
clang/test/SemaCUDA/global-initializers.cu
Removed:
################################################################################
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 13aa06aa97399..090030ea82503 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -776,39 +776,63 @@ 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 &&
+ !VD->isInvalidDecl()) {
+ 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);
+ }
+ }
+ void VisitCallExpr(const CallExpr *CE) {
+ const FunctionDecl *InitFn = CE->getDirectCallee();
+ if (InitFn)
+ CheckForWrongSidedCall(InitFn);
+ Base::VisitCallExpr(CE);
+ }
+
+ void VisitCXXConstructExpr(const CXXConstructExpr *CE) {
+ 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->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..cba5c539614c7 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,69 @@ __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 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 {
+ 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}}
+ }
+
+ 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}}
+
+ 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