https://github.com/Artem-B created
https://github.com/llvm/llvm-project/pull/142857
The variables have implicit host-side shadow instances and explicit address
space attribute breaks them on the host.
>From e2e8da0271ae11711dbd54f6e8d9ff498f3226d4 Mon Sep 17 00:00:00 2001
From: Artem Belevich <t...@google.com>
Date: Tue, 27 May 2025 17:09:01 -0700
Subject: [PATCH] [CUDA] Disallow use of address_space(N) on CUDA device
variables.
The variables have implicit host-side shadow instances and explicit address
space attribute breaks them on the host.
---
.../include/clang/Basic/DiagnosticSemaKinds.td | 2 ++
clang/lib/Sema/SemaCUDA.cpp | 17 +++++++++++------
clang/test/SemaCUDA/bad-attributes.cu | 8 ++++++++
3 files changed, 21 insertions(+), 6 deletions(-)
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}}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
