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@+3 {{'value_func' declared here}} + // expected-note@+2 9{{'value_func' declared here}} + // expected-note@+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@-1 {{no matching function for call to 'value_func'}} + + // diagnosed via SemaCUDA::checkAllowedInitializer + constexpr static unsigned var1 = ns1::value_func(); + // host-error@-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@-1 {{reference to __device__ function 'value_func' in global initializer}} + static int var3 {ns1::value_func()}; + // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} + + struct A { + unsigned b; + }; + A b{ns1::value_func()}; + // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} + + int foo(int); + int nested = foo(ns1::value_func()); + // expected-error@-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@-1 {{reference to __device__ function 'value_func' in __host__ function}} + // device-error@-2 {{reference to __device__ function 'value_func' in global initializer}} + } + + struct DefInit { + unsigned a = 1 + value_func(); + }; + DefInit testDefInit; + // expected-error@-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@-1 {{reference to __device__ function 'value_func' in global initializer}} + + unsigned twotimes = ns1::value_func() + ns1::another_value_func(); + // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} + + unsigned twotimes1 = ns1::value_func() + ns1::value_func(); + // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} + +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
