From 73b22935a7a863679021598db6a45fcfb62cd321 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Fri, 11 Feb 2022 23:07:46 -0500 Subject: [PATCH] [CUDA][HIP] Do not promote constexpr var with non-constant initializer constexpr var may be initialized with address of non-const variable. In this case the initializer is not constant in device compilation. This has been handled for const vars but not for constexpr vars. This patch makes handling of const var and constexpr var consistent. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D119615 Fixes: https://github.com/llvm/llvm-project/issues/53780 --- clang/lib/Sema/SemaCUDA.cpp | 14 +++-- clang/test/SemaCUDA/constexpr-var.cu | 105 +++++++++++++++++++++++++++++++++++ 2 files changed, 113 insertions(+), 6 deletions(-) create mode 100644 clang/test/SemaCUDA/constexpr-var.cu diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp index efa3855..e4e34d6 100644 --- a/clang/lib/Sema/SemaCUDA.cpp +++ b/clang/lib/Sema/SemaCUDA.cpp @@ -145,9 +145,11 @@ Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D, Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) { if (Var->hasAttr()) return CVT_Unified; - if (Var->isConstexpr() && !hasExplicitAttr(Var)) - return CVT_Both; - if (Var->getType().isConstQualified() && Var->hasAttr() && + // Only constexpr and const variabless with implicit constant attribute + // are emitted on both sides. Such variables are promoted to device side + // only if they have static constant intializers on device side. + if ((Var->isConstexpr() || Var->getType().isConstQualified()) && + Var->hasAttr() && !hasExplicitAttr(Var)) return CVT_Both; if (Var->hasAttr() || Var->hasAttr() || @@ -718,9 +720,9 @@ void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) { !VD->hasAttr() && !VD->hasAttr() && (VD->isFileVarDecl() || VD->isStaticDataMember()) && !IsDependentVar(VD) && - (VD->isConstexpr() || (VD->getType().isConstQualified() && - HasAllowedCUDADeviceStaticInitializer( - *this, VD, CICK_DeviceOrConstant)))) { + ((VD->isConstexpr() || VD->getType().isConstQualified()) && + HasAllowedCUDADeviceStaticInitializer(*this, VD, + CICK_DeviceOrConstant))) { VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext())); } } diff --git a/clang/test/SemaCUDA/constexpr-var.cu b/clang/test/SemaCUDA/constexpr-var.cu new file mode 100644 index 0000000..a028ba8f --- /dev/null +++ b/clang/test/SemaCUDA/constexpr-var.cu @@ -0,0 +1,105 @@ +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \ +// RUN: -fsyntax-only -verify +// RUN: %clang_cc1 -triple x86_64 -x hip %s \ +// RUN: -fsyntax-only -verify=host + +// host-no-diagnostics + +#include "Inputs/cuda.h" + +// Test constexpr var initialized with address of a const var. +// Both are promoted to device side. + +namespace Test1 { +const int a = 1; + +struct B { + static constexpr const int *p = &a; + __device__ static constexpr const int *const p2 = &a; +}; + +// Const variable 'a' is treated as __constant__ on device side, +// therefore its address can be used as initializer for another +// device variable. + +__device__ void f() { + int y = a; + constexpr const int *x = B::p; + constexpr const int *z = B::p2; +} +} + +// Test constexpr var initialized with address of a non-cost var. +// Neither is promoted to device side. + +namespace Test2 { +int a = 1; +// expected-note@-1{{host variable declared here}} + +struct B { + static constexpr int *const p = &a; + // expected-note@-1{{const variable cannot be emitted on device side due to dynamic initialization}} +}; + +__device__ void f() { + int y = a; + // expected-error@-1{{reference to __host__ variable 'a' in __device__ function}} + const int *const *x = &B::p; + // expected-error@-1{{reference to __host__ variable 'p' in __device__ function}} + // ToDo: use of non-promotable constexpr variable in device compilation should be treated as + // ODR-use and diagnosed. + const int *const z = B::p; +} +} + +// Test constexpr device var initialized with address of a non-const host var, __shared var, +// __managed__ var, __device__ var, __constant__ var, texture var, surface var. + +namespace Test3 { +struct textureReference { + int desc; +}; + +enum ReadMode { + ElementType = 0, + NormalizedFloat = 1 +}; + +template +struct __attribute__((device_builtin_texture_type)) texture : public textureReference { +}; + +struct surfaceReference { + int desc; +}; + +template +struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference { +}; + +// Partial specialization over `void`. +template +struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference { +}; + +texture tex; +surface surf; + +int a = 1; +__shared__ int b; +__managed__ int c = 1; +__device__ int d = 1; +__constant__ int e = 1; +struct B { + __device__ static constexpr int *const p1 = &a; + // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}} + __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}} + __device__ static constexpr int *const p4 = &d; + __device__ static constexpr int *const p5 = &e; + __device__ static constexpr texture *const p6 = &tex; + __device__ static constexpr surface *const p7 = &surf; +}; +} -- 2.7.4