https://github.com/Fznamznon updated https://github.com/llvm/llvm-project/pull/217047
>From ccfa00602d3348cef26dc9b1e285a42426b564a8 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Tue, 18 Aug 2026 09:57:12 -0500 Subject: [PATCH 1/2] [clang][HIP] Do not treat address of managed variable as a constant expression According to CUDA programming guilde, the address of a __managed__ variable is not a constant expression, so it should not be accepted where a constant expression is expected, i.e. in NTTP, constexpr variable initializers etc. Right now because addresses of managed variables are assumed to be constexpr, crashes in clang's codegen happen during replacement of uses of managed variables with loads from transformed managed variables. It is not expected that a use of a managed variable can be `llvm::ConstantExpr` which clang creates due to assumption that address of a managed variable is a constant expression. Fixes https://github.com/llvm/llvm-project/issues/198079 --- clang/lib/AST/ExprConstant.cpp | 7 +++++-- clang/test/CodeGenCUDA/managed-var.cu | 20 ++++++++++++++++++++ clang/test/SemaCUDA/const-var.cu | 3 +-- clang/test/SemaCUDA/constexpr-var.cu | 9 ++++----- clang/test/SemaCUDA/managed-var.cu | 22 ++++++++++++++++++++++ 5 files changed, 52 insertions(+), 9 deletions(-) diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index 05f981f671e62..feb141a7d141c 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -2313,6 +2313,10 @@ static bool CheckLValueConstantExpression(EvalInfo &Info, SourceLocation Loc, !Var->isStaticLocal()) return false; + // Address of a managed variable is never a constant expression. + if (Info.getLangOpts().CUDA && Var->hasAttr<HIPManagedAttr>()) + return false; + // In CUDA/HIP device compilation, only device side variables have // constant addresses. if (Info.getLangOpts().CUDA && Info.getLangOpts().CUDAIsDevice && @@ -2320,8 +2324,7 @@ static bool CheckLValueConstantExpression(EvalInfo &Info, SourceLocation Loc, if ((!Var->hasAttr<CUDADeviceAttr>() && !Var->hasAttr<CUDAConstantAttr>() && !Var->getType()->isCUDADeviceBuiltinSurfaceType() && - !Var->getType()->isCUDADeviceBuiltinTextureType()) || - Var->hasAttr<HIPManagedAttr>()) + !Var->getType()->isCUDADeviceBuiltinTextureType())) return false; } } diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu index c2ec9d433e40a..081099c0b6dcb 100644 --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -161,6 +161,26 @@ __device__ __host__ int load4() { return ex; } +namespace gh198079 { +__managed__ int x = 0; +struct S { + int *p; +}; +__attribute__((device)) void f() { + S s{&x}; +} +// DEV-LABEL: define {{.*}}@{{.*}}gh198079{{.*}}f{{.*}}() +// DEV: %p = getelementptr inbounds nuw %"struct.gh198079::S", ptr %s.ascast, i32 0, i32 0 +// DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @_ZN8gh1980791xE, align 4 +// DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr +// DEV: store ptr %0, ptr %p, align 8 +int *hostglob = &x; + +// HOST-LABEL: define internal void @__cxx_global_var_init() +// HOST: %ld.managed = load ptr, ptr @_ZN8gh1980791xE +// HOST: store ptr %ld.managed, ptr @_ZN8gh1980798hostglobE +} + // HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @x, ptr @x.managed, ptr @[[DEVNAMEX]], i64 4, i32 4) // HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @_ZL2sx, ptr @_ZL2sx.managed, ptr @[[DEVNAMESX]] // HOST-NOT: __hipRegisterManagedVar({{.*}}, ptr @ex, ptr @ex.managed diff --git a/clang/test/SemaCUDA/const-var.cu b/clang/test/SemaCUDA/const-var.cu index 22317d0f0ee0c..3ed5273f430f0 100644 --- a/clang/test/SemaCUDA/const-var.cu +++ b/clang/test/SemaCUDA/const-var.cu @@ -3,8 +3,6 @@ // RUN: %clang_cc1 -triple x86_64 -x hip %s \ // RUN: -fsyntax-only -verify=host -// host-no-diagnostics - #include "Inputs/cuda.h" // Test const var initialized with address of a const var. @@ -104,6 +102,7 @@ __device__ int *const B::p2 = &b; // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} __device__ int *const B::p3 = &c; // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} +// host-error@-2{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} __device__ int *const B::p4 = &d; __device__ int *const B::p5 = &e; __device__ texture<float, 2, ElementType> *const B::p6 = &tex; diff --git a/clang/test/SemaCUDA/constexpr-var.cu b/clang/test/SemaCUDA/constexpr-var.cu index 5bb6cc1208c98..0df852c7adbeb 100644 --- a/clang/test/SemaCUDA/constexpr-var.cu +++ b/clang/test/SemaCUDA/constexpr-var.cu @@ -1,9 +1,7 @@ // RUN: %clang_cc1 -triple amdgpu-amd-amdhsa -fcuda-is-device -x hip %s \ -// RUN: -fsyntax-only -verify +// RUN: -fsyntax-only -verify=expected,both // RUN: %clang_cc1 -triple x86_64 -x hip %s \ -// RUN: -fsyntax-only -verify=host - -// host-no-diagnostics +// RUN: -fsyntax-only -verify=host,both #include "Inputs/cuda.h" @@ -96,7 +94,8 @@ struct B { __device__ static constexpr int *const p2 = &b; // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} __device__ static constexpr int *const p3 = &c; - // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + // both-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + // both-error@-2{{constexpr variable 'p3' must be initialized by a constant expression}} __device__ static constexpr int *const p4 = &d; __device__ static constexpr int *const p5 = &e; __device__ static constexpr texture<float, 2, ElementType> *const p6 = &tex; diff --git a/clang/test/SemaCUDA/managed-var.cu b/clang/test/SemaCUDA/managed-var.cu index 3f699b79a0437..68613460a0a9f 100644 --- a/clang/test/SemaCUDA/managed-var.cu +++ b/clang/test/SemaCUDA/managed-var.cu @@ -52,3 +52,25 @@ typedef __managed__ int managed_int; __managed__ A a; // expected-error@-1 {{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + +namespace gh198079 { +__managed__ int x = 0; +template <int *P> int *get() { return P; } // expected-note {{ ignored: non-type template argument is not a constant expression}} + +void foo() { + static constexpr auto a = &x; // expected-error {{constexpr variable 'a' must be initialized by a constant expression}} + get<&x>(); // expected-error {{no matching function for call to 'get'}} + +} + +template <int *PP> +class boop { +public: + static constexpr auto B = PP; +}; + +__device__ void bar() { + static constexpr auto A = boop<&x>::B; // expected-error {{non-type template argument is not a constant expression}} +} + +} >From 1ca612b5a2737173a564e068cda3f01c895667dd Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Wed, 19 Aug 2026 07:46:05 -0500 Subject: [PATCH 2/2] Fix the test --- clang/test/CodeGenCUDA/managed-var.cu | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu index 081099c0b6dcb..96b61d1313d2f 100644 --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -163,19 +163,22 @@ __device__ __host__ int load4() { namespace gh198079 { __managed__ int x = 0; + struct S { - int *p; + int *p; }; -__attribute__((device)) void f() { + +__device__ __host__ void f() { S s{&x}; } -// DEV-LABEL: define {{.*}}@{{.*}}gh198079{{.*}}f{{.*}}() -// DEV: %p = getelementptr inbounds nuw %"struct.gh198079::S", ptr %s.ascast, i32 0, i32 0 +// COMMON-LABEL: define {{.*}}@{{.*}}gh198079{{.*}}f{{.*}}() // DEV: %ld.managed = load ptr addrspace(1), ptr addrspace(1) @_ZN8gh1980791xE, align 4 // DEV: %0 = addrspacecast ptr addrspace(1) %ld.managed to ptr -// DEV: store ptr %0, ptr %p, align 8 -int *hostglob = &x; +// DEV: store ptr %0, ptr %p +// HOST: %ld.managed = load ptr, ptr @_ZN8gh1980791xE, align 4 +// HOST: store ptr %ld.managed, ptr %p +int *hostglob = &x; // HOST-LABEL: define internal void @__cxx_global_var_init() // HOST: %ld.managed = load ptr, ptr @_ZN8gh1980791xE // HOST: store ptr %ld.managed, ptr @_ZN8gh1980798hostglobE _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
