Author: Mariya Podchishchaeva Date: 2026-08-25T09:05:32+02:00 New Revision: dad29286872ab2f59224042c3acb3c3a8c0a904b
URL: https://github.com/llvm/llvm-project/commit/dad29286872ab2f59224042c3acb3c3a8c0a904b DIFF: https://github.com/llvm/llvm-project/commit/dad29286872ab2f59224042c3acb3c3a8c0a904b.diff LOG: [clang][HIP] Do not treat address of managed variable as a constant expression (#217047) According to CUDA programming guide, 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. 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 Assisted-by: claude in documentation writing Added: Modified: clang/docs/HIPSupport.md clang/include/clang/Basic/AttrDocs.td clang/include/clang/Basic/DiagnosticSemaKinds.td clang/lib/AST/ExprConstant.cpp clang/lib/Sema/SemaCUDA.cpp clang/test/CodeGenCUDA/managed-var.cu clang/test/SemaCUDA/const-var.cu clang/test/SemaCUDA/constexpr-var.cu clang/test/SemaCUDA/managed-var.cu Removed: ################################################################################ diff --git a/clang/docs/HIPSupport.md b/clang/docs/HIPSupport.md index 8980a27c55474..892d411520afe 100644 --- a/clang/docs/HIPSupport.md +++ b/clang/docs/HIPSupport.md @@ -407,6 +407,64 @@ __host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv"))); __host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf"))); ``` +## Managed Variables + +Clang currently implements the following diagnostics involving +`__managed__` variables: + +- 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: + + ```c++ + __managed__ int x; + + void foo() { + static constexpr auto a = &x; // error: constexpr variable 'a' must be + // initialized by a constant expression + } + ``` + +- Managed variables 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. + ``` + +A managed variable is accessible from both host and device and its host and +device initializers must be consistent; otherwise the behavior is undefined. +A managed variable is emitted as an undefined global symbol in the device binary +and is registered with the HIP runtime by `__hipRegisterManagedVar` function +call during device module loading. Clang replaces device accesses to +a managed variable with loads from a pointer to a chunk of managed memory +allocated by the HIP runtime. + ## C++17 Class Template Argument Deduction (CTAD) Support Clang supports C++17 Class Template Argument Deduction (CTAD) in both host and diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 3052dd6c77ab1..5ec2a68bc038b 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -7922,7 +7922,7 @@ def HIPManagedAttrDocs : Documentation { let Content = [{ The `__managed__` attribute can be applied to a global variable declaration in HIP. A managed variable is emitted as an undefined global symbol in the device binary and is -registered by `__hipRegisterManagedVariable` in init functions. The HIP runtime allocates +registered by `__hipRegisterManagedVar` in init functions. The HIP runtime allocates managed memory and uses it to define the symbol when loading the device binary. A managed variable can be accessed in both device and host code. }]; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3a910c9c3f2b9..384ff992de534 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/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index d55749100658f..e5a1d1fe56a56 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -2314,6 +2314,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 && @@ -2321,8 +2325,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/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index a2a088ab7c3ab..78cfa2c1a8662 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/EvaluatedExprVisitor.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 : ConstEvaluatedExprVisitor<GlobVarInitChecker> { + using Base = ConstEvaluatedExprVisitor<GlobVarInitChecker>; + SemaCUDA &SCRef; + SourceLocation InitLoc; + + GlobVarInitChecker(SemaCUDA &S, SourceLocation L) + : 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); + } + } + }; + GlobVarInitChecker Checker(*this, VD->getLocation()); + Checker.Visit(Init); } } diff --git a/clang/test/CodeGenCUDA/managed-var.cu b/clang/test/CodeGenCUDA/managed-var.cu index c2ec9d433e40a..5a087cc01699b 100644 --- a/clang/test/CodeGenCUDA/managed-var.cu +++ b/clang/test/CodeGenCUDA/managed-var.cu @@ -161,6 +161,24 @@ __device__ __host__ int load4() { return ex; } +namespace gh198079 { +__managed__ int x = 0; + +struct S { + int *p; +}; + +__device__ __host__ void f() { + S s{&x}; +} +// 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 +// HOST: %ld.managed = load ptr, ptr @_ZN8gh1980791xE, align 4 +// HOST: store ptr %ld.managed, ptr %p +} + // 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..5e05c0d25f6e3 100644 --- a/clang/test/SemaCUDA/managed-var.cu +++ b/clang/test/SemaCUDA/managed-var.cu @@ -52,3 +52,55 @@ 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@-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'}} + +} + +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}} +} +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}} + +int ok = sizeof(x); +auto reader = [] { return x; }; +} // namespace gh198079 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
