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/6] [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/6] 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 >From 2e29cfd87f614ef39723e35b79180a490b45fc08 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Wed, 19 Aug 2026 09:24:25 -0500 Subject: [PATCH 3/6] Diagnose uses of managed variables in glob var inits --- .../clang/Basic/DiagnosticSemaKinds.td | 5 +++ clang/lib/Sema/SemaCUDA.cpp | 19 ++++++++++++ clang/test/CodeGenCUDA/managed-var.cu | 5 --- clang/test/SemaCUDA/managed-var.cu | 31 +++++++++++++++++-- 4 files changed, 53 insertions(+), 7 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 22e65d3205808..472eab031283c 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9844,6 +9844,11 @@ def err_cuda_host_shared : Error< "%select{__device__|__global__|__host__|__host__ __device__}0 functions">; def err_cuda_nonstatic_constdev: Error<"__constant__, __device__, and " "__managed__ are not allowed on non-static local variables">; +def err_cuda_invalid_use_of_managedvar: Error<"invalid use of a __managed__ " + "variable">; +def note_cuda_managed_var_in_glob_init : Note<"__managed__ variable shall not " + "be used in initialization of an object with static or thread local " + "storage duration">; def err_cuda_address_space_gpuvar: Error<"__constant__, __device__, and " "__shared__ variables must use default address space">; def err_cuda_grid_constant_not_allowed : Error< diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index a2a088ab7c3ab..886682df5fb8e 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -13,6 +13,7 @@ #include "clang/Sema/SemaCUDA.h" #include "clang/AST/ASTContext.h" #include "clang/AST/Decl.h" +#include "clang/AST/DynamicRecursiveASTVisitor.h" #include "clang/AST/ExprCXX.h" #include "clang/Basic/Cuda.h" #include "clang/Basic/TargetInfo.h" @@ -789,6 +790,24 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { VD->setInvalidDecl(); } } + struct GlobVarInitChecker : DynamicRecursiveASTVisitor { + SemaCUDA &SCRef; + SourceLocation InitLoc; + + GlobVarInitChecker(SemaCUDA &S, SourceLocation L) + : SCRef(S), InitLoc(L) {} + bool VisitDeclRefExpr(DeclRefExpr *DRE) override { + 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); + } + return true; + } + }; + GlobVarInitChecker Checker(*this, VD->getLocation()); + Checker.TraverseStmt(VD->getInit()); } } diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu index 96b61d1313d2f..5a087cc01699b 100644 --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -177,11 +177,6 @@ __device__ __host__ void f() { // 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 } // HOST-DAG: __hipRegisterManagedVar({{.*}}, ptr @x, ptr @x.managed, ptr @[[DEVNAMEX]], i64 4, i32 4) diff --git a/clang/test/SemaCUDA/managed-var.cu b/clang/test/SemaCUDA/managed-var.cu index 68613460a0a9f..13df82ef4f25b 100644 --- a/clang/test/SemaCUDA/managed-var.cu +++ b/clang/test/SemaCUDA/managed-var.cu @@ -58,7 +58,11 @@ __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}} + static constexpr auto a = &x; + // expected-error@-1 {{constexpr variable 'a' must be initialized by a constant expression}} + // expected-error@-2 {{invalid use of a __managed__ variable}} + // expected-note@-3 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} + get<&x>(); // expected-error {{no matching function for call to 'get'}} } @@ -72,5 +76,28 @@ public: __device__ void bar() { static constexpr auto A = boop<&x>::B; // expected-error {{non-type template argument is not a constant expression}} } - +int *hostglob = &x; +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} +int hostglob1 = x; +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} +size_t hostglob2 = (size_t)&x; +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} +struct Test { + Test(int) {} +}; +Test hostglob3{x}; +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} + +int foo(int); +Test hostglob4(foo(x)); +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} + +thread_local int hostglob5 = x; +// expected-error@-1 {{invalid use of a __managed__ variable}} +// expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} } >From 38cfab2b7156d0da72cc747e4755132d66e506be Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Thu, 20 Aug 2026 08:10:17 -0500 Subject: [PATCH 4/6] Add documentation --- clang/docs/HIPSupport.md | 57 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 57 insertions(+) diff --git a/clang/docs/HIPSupport.md b/clang/docs/HIPSupport.md index 8980a27c55474..43d12f1d8b54d 100644 --- a/clang/docs/HIPSupport.md +++ b/clang/docs/HIPSupport.md @@ -407,6 +407,63 @@ __host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv"))); __host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf"))); ``` +## Managed Variables + +The `__managed__` attribute can be applied to a global variable to indicate that +the variable's memory is accessible by both host and device. +A managed variable is emitted as an undefined global symbol in the device binary +and is registered with the HIP runtime by `__hipRegisterManagedVariable` in the +module init functions. + +### Restrictions +`__managed__` variables have the following restrictions: + +- No dynamic initialization. Only constant initialization is permitted: + + ```c++ + struct A { + int a; + A() { a = 1; } + }; + + __managed__ A a; // error: dynamic initialization is not supported + + ``` + +- The address of a managed variable is not a constant expression, because + the runtime allocates the storage and defines the symbol at load time: + + ```c++ + __managed__ int x; + + void foo() { + static constexpr auto a = &x; // error: constexpr variable 'a' must be + // initialized by a constant expression + } + ``` + +- Shall not be used in the initializer of a static/thread-local object, + because that initialization may run before the HIP runtime has registered the + managed variable, leaving its address null. + Clang can diagnose direct uses of a managed variable within an initializer; + however, it will not follow function calls, constructors, + destructors, function pointers, or definitions in other translation units, + i.e. even if a managed variable is used there, an error won't be emitted. + Accessing a managed variable during static or thread-local initialization or + destruction is still undefined behavior, even if Clang does not emit a + diagnostic. Example: + + ```c++ + __managed__ int x; + + int *hostglob = &x; // error: invalid use of a __managed__ variable + + int *getX() { + return &x; + } + int *p = getX(); // No error emitted. + ``` + ## C++17 Class Template Argument Deduction (CTAD) Support Clang supports C++17 Class Template Argument Deduction (CTAD) in both host and >From cef9926132b5622ad1d2462e608a177ff748ff0d Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Thu, 20 Aug 2026 08:12:40 -0500 Subject: [PATCH 5/6] Switch to ConstEvaluatedExprVisitor --- clang/lib/Sema/SemaCUDA.cpp | 12 ++++++------ clang/test/SemaCUDA/managed-var.cu | 4 +++- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index 886682df5fb8e..78cfa2c1a8662 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -13,7 +13,7 @@ #include "clang/Sema/SemaCUDA.h" #include "clang/AST/ASTContext.h" #include "clang/AST/Decl.h" -#include "clang/AST/DynamicRecursiveASTVisitor.h" +#include "clang/AST/EvaluatedExprVisitor.h" #include "clang/AST/ExprCXX.h" #include "clang/Basic/Cuda.h" #include "clang/Basic/TargetInfo.h" @@ -790,24 +790,24 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) { VD->setInvalidDecl(); } } - struct GlobVarInitChecker : DynamicRecursiveASTVisitor { + struct GlobVarInitChecker : ConstEvaluatedExprVisitor<GlobVarInitChecker> { + using Base = ConstEvaluatedExprVisitor<GlobVarInitChecker>; SemaCUDA &SCRef; SourceLocation InitLoc; GlobVarInitChecker(SemaCUDA &S, SourceLocation L) - : SCRef(S), InitLoc(L) {} - bool VisitDeclRefExpr(DeclRefExpr *DRE) override { + : Base(S.getASTContext()), SCRef(S), InitLoc(L) {} + 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); } - return true; } }; GlobVarInitChecker Checker(*this, VD->getLocation()); - Checker.TraverseStmt(VD->getInit()); + Checker.Visit(Init); } } diff --git a/clang/test/SemaCUDA/managed-var.cu b/clang/test/SemaCUDA/managed-var.cu index 13df82ef4f25b..bb4bedd8ed9ce 100644 --- a/clang/test/SemaCUDA/managed-var.cu +++ b/clang/test/SemaCUDA/managed-var.cu @@ -100,4 +100,6 @@ Test hostglob4(foo(x)); thread_local int hostglob5 = x; // expected-error@-1 {{invalid use of a __managed__ variable}} // expected-note@-2 {{__managed__ variable shall not be used in initialization of an object with static or thread local storage duration}} -} + +int ok = sizeof(x); +} // namespace gh198079 >From d656830295f0db875edd8ba5d6d26dee071c3d79 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva <[email protected]> Date: Thu, 20 Aug 2026 08:18:07 -0500 Subject: [PATCH 6/6] Shorten doc a bit --- clang/docs/HIPSupport.md | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/docs/HIPSupport.md b/clang/docs/HIPSupport.md index 43d12f1d8b54d..f1fb77b1bfdf5 100644 --- a/clang/docs/HIPSupport.md +++ b/clang/docs/HIPSupport.md @@ -430,8 +430,7 @@ module init functions. ``` -- The address of a managed variable is not a constant expression, because - the runtime allocates the storage and defines the symbol at load time: +- The address of a managed variable is not a constant expression: ```c++ __managed__ int x; _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
