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

[Artem-B]

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

[llvmbot]

@llvm/pr-subscribers-clang

Author: Artem Belevich (Artem-B)

Changes

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

---

Full diff: https://github.com/llvm/llvm-project/pull/142857.diff

3 Files Affected:

• (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+2)
• (modified) clang/lib/Sema/SemaCUDA.cpp (+11-6)
• (modified) clang/test/SemaCUDA/bad-attributes.cu (+8)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 78b36ceb88125..3862454ed2ee9 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9388,6 +9388,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<
diff --git a/clang/lib/Sema/SemaCUDA.cpp b/clang/lib/Sema/SemaCUDA.cpp
index 176d9322d3d2a..67e7a2ed61db9 100644
--- a/clang/lib/Sema/SemaCUDA.cpp
+++ b/clang/lib/Sema/SemaCUDA.cpp
@@ -322,7 +322,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) {
@@ -505,7 +505,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;
@@ -680,16 +679,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))
diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
index 8ac6fa1ab5c2d..acb2f38a75282 100644
--- a/clang/test/SemaCUDA/bad-attributes.cu
+++ b/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}}

[alexfh]

Looks good. Thanks for adding this early detection of the incorrect CUDA code that would otherwise lead to violated invariants further down the road.