caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
The variables have implicit host-side shadow instances and explicit
address space attribute breaks them on the host.
This commit is contained in:
Artem Belevich
2025-06-09 10:39:04 -07:00
committed by GitHub
parent 649020c680
commit 59ef2c34a1
3 changed files with 21 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
clang/include/clang/Basic/DiagnosticSemaKinds.td
View File

@@ -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<

17
clang/lib/Sema/SemaCUDA.cpp
View File

@@ -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))

8
clang/test/SemaCUDA/bad-attributes.cu
View File

@@ -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}}