[CUDA] Disallow use of address_space(N) on CUDA device variables. (llvm#142857)

The variables have implicit host-side shadow instances and explicit
address space attribute breaks them on the host.

Author: Artem-B

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9432,6 +9432,8 @@ 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_address_space_gpuvar: Error<"__constant__, __device__, and "
+    "__shared__ variables must use default address space">;
 def err_cuda_grid_constant_not_allowed : Error<
   "__grid_constant__ is only allowed on const-qualified kernel parameters">;
 def err_cuda_ovl_target : Error<

clang/lib/Sema/SemaCUDA.cpp

@@ -321,7 +321,7 @@ void SemaCUDA::EraseUnwantedMatches(
   if (Matches.size() <= 1)
     return;
 
-  using Pair = std::pair<DeclAccessPair, FunctionDecl*>;
+  using Pair = std::pair<DeclAccessPair, FunctionDecl *>;
 
   // Gets the CUDA function preference for a call from Caller to Match.
   auto GetCFP = [&](const Pair &Match) {
@@ -504,7 +504,6 @@ bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
     }
   }
 
-
   // If no target was inferred, mark this member as __host__ __device__;
   // it's the least restrictive option that can be invoked from any target.
   bool NeedsH = true, NeedsD = true;
@@ -679,16 +678,22 @@ void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
       FD && FD->isDependentContext())
     return;
 
+  bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
+  bool IsDeviceOrConstantVar =
+      !IsSharedVar &&
+      (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
+  if ((IsSharedVar || IsDeviceOrConstantVar) &&
+      VD->getType().getQualifiers().getAddressSpace() != LangAS::Default) {
+    Diag(VD->getLocation(), diag::err_cuda_address_space_gpuvar);
+    VD->setInvalidDecl();
+    return;
+  }
   // Do not check dependent variables since the ctor/dtor/initializer are not
   // determined. Do it after instantiation.
   if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() ||
       IsDependentVar(VD))
     return;
   const Expr *Init = VD->getInit();
-  bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
-  bool IsDeviceOrConstantVar =
-      !IsSharedVar &&
-      (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
   if (IsDeviceOrConstantVar || IsSharedVar) {
     if (HasAllowedCUDADeviceStaticInitializer(
             *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant))

clang/test/SemaCUDA/bad-attributes.cu

@@ -50,6 +50,14 @@ __global__ __device__ void z11();  // expected-error {{attributes are not compat
 __global__ __host__ void z12();  // expected-error {{attributes are not compatible}}
 // expected-note@-1 {{conflicting attribute is here}}
 
+// Make sure GPU-side variables do not allow __attribute((address_space(N)))
+// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}
+__shared__ __attribute__((address_space(999))) int as_s;
+// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}
+__device__ __attribute__((address_space(999))) int as_d;
+// expected-error@+1 {{__constant__, __device__, and __shared__ variables must use default address space}}
+__constant__ __attribute__((address_space(999))) int as_c;
+
 struct S {
   __global__ void foo() {};  // expected-error {{must be a free function or static member function}}
   __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}}