https://github.com/Fznamznon updated https://github.com/llvm/llvm-project/pull/233925
>From fb66044df94f2d1cad3fd46b79cfeec086bc58f6 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Tue, 20 Oct 2026 12:37:33 -0500 Subject: [PATCH 1/3] [CUDA] Fix diagnosing of static variable initializers in device functions 4735d351182f874d61690f6fc2d2e3b154914ad5 extended diagnostics for __device__ functions used from host global variable initializers which caused errors to be emitted more frequently and revealed that the case that was intended only for host variables also false positively errors on static variables inside of device functions that are initialized using other device functions. So, before checking make sure the variable is host variable. --- clang/lib/Sema/SemaCUDA.cpp | 3 ++- clang/test/SemaCUDA/global-initializers.cu | 8 ++++++++ 2 files changed, 10 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index fbdd7b0a40ae72..de979e769e6704 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -766,6 +766,7 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { IsDependentVar(VD)) return; const Expr *Init = VD->getInit(); + if (IsDeviceOrConstantVar || IsSharedVar) { if (HasAllowedCUDADeviceStaticInitializer( *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant)) @@ -774,7 +775,7 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init) << Init->getSourceRange(); VD->setInvalidDecl(); - } else { + } else if (IdentifyTarget(VD) == CVT_Host) { // This is a host-side global variable. Check that the initializer is // callable from the host side. diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu index cba5c539614c7e..0c4adfde52d79b 100644 --- a/clang/test/SemaCUDA/global-initializers.cu +++ b/clang/test/SemaCUDA/global-initializers.cu @@ -135,4 +135,12 @@ namespace ns2 { unsigned twotimes1 = ns1::value_func() + ns1::value_func(); // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} + __device__ constexpr int devicefoo() { return 1; } + + __device__ void static testfps() { + static constexpr int b = value_func(); + static int b1 = devicefoo(); + [](){ static int b2 = devicefoo(); }(); + } + } >From 2c1e78dc8f2905dc19a45328435d884aacf6019e Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Tue, 20 Oct 2026 12:45:23 -0500 Subject: [PATCH 2/3] Remove empty line --- 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 de979e769e6704..e6a92137395de4 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -766,7 +766,6 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { IsDependentVar(VD)) return; const Expr *Init = VD->getInit(); - if (IsDeviceOrConstantVar || IsSharedVar) { if (HasAllowedCUDADeviceStaticInitializer( *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant)) >From 8df4f84d5f82bcd2ee78772a34c7274af7d93f1c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Wed, 21 Oct 2026 06:23:21 -0500 Subject: [PATCH 3/3] Preserve host check for HD functions, add tests --- clang/lib/Sema/SemaCUDA.cpp | 4 +- clang/test/SemaCUDA/global-initializers.cu | 44 ++++++++++++++++++++-- 2 files changed, 44 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index e6a92137395de4..ffad8097f59a5a 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -766,6 +766,7 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { IsDependentVar(VD)) return; const Expr *Init = VD->getInit(); + SemaCUDA::CUDAVariableTarget VarTarget = IdentifyTarget(VD); if (IsDeviceOrConstantVar || IsSharedVar) { if (HasAllowedCUDADeviceStaticInitializer( *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant)) @@ -774,7 +775,8 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init) << Init->getSourceRange(); VD->setInvalidDecl(); - } else if (IdentifyTarget(VD) == CVT_Host) { + } else if (VarTarget == CVT_Host || + (!getLangOpts().CUDAIsDevice && VarTarget == CVT_Both)) { // This is a host-side global variable. Check that the initializer is // callable from the host side. diff --git a/clang/test/SemaCUDA/global-initializers.cu b/clang/test/SemaCUDA/global-initializers.cu index 0c4adfde52d79b..5ddf8749239c3d 100644 --- a/clang/test/SemaCUDA/global-initializers.cu +++ b/clang/test/SemaCUDA/global-initializers.cu @@ -136,11 +136,49 @@ namespace ns2 { // expected-error@-1 {{reference to __device__ function 'value_func' in global initializer}} __device__ constexpr int devicefoo() { return 1; } + // host-note@-1 6{{declared here}} - __device__ void static testfps() { - static constexpr int b = value_func(); + __device__ void static test_statics_in_device() { + static constexpr int b = devicefoo(); static int b1 = devicefoo(); - [](){ static int b2 = devicefoo(); }(); + [](){ // lambdas are automatically HD + static int b2 = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + static constexpr int b3 = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + }(); + + + template<class T> + static constexpr T TVar = devicefoo; + + int a = TVar<int> + 1; + } + __host__ __device__ void static test_statics_in_hd() { + static constexpr int b = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + static int b1 = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + [](){ + static int b2 = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + static constexpr int b3 = devicefoo(); + // host-error@-1 {{reference to __device__ function 'devicefoo' in global initializer}} + }(); } + template <typename T> __device__ void templatedevice() { + static constexpr T b = devicefoo(); + } + template <> + __device__ void templatedevice<int>(); + template <typename T> __device__ __host__ void templatehd() { + static constexpr T b = devicefoo(); + } + // FIXME: probably should be diagnosed for host + template <> + __device__ __host__ void templatehd<int>(); + + + } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
