[clang] [llvm] [clang][CUDA] Allow `extern __shared__` on non-array types (PR #183879)
Yuxuan Chen via cfe-commits
cfe-commits at lists.llvm.org
Sun Oct 25 23:28:51 PDT 2026
https://github.com/yuxuanchen1997 updated https://github.com/llvm/llvm-project/pull/183879
>From c53bb5e5ad81b6963c3da2ba33629868a0a62798 Mon Sep 17 00:00:00 2001
From: Yuxuan Chen <ych at meta.com>
Date: Fri, 27 Feb 2026 15:57:24 -0800
Subject: [PATCH] [clang][CUDA] Allow extern __shared__ on complete types
Accept extern __shared__ scalars, structs, and fixed-size arrays for NVCC
compatibility. Replace the non-RDC error with a default-ignored warning
under -Wcuda-extern-shared and retain CUDASharedAttr so device references
are classified correctly.
In non-RDC mode, nonzero-sized complete types use fixed shared storage;
unsized arrays use dynamic shared memory. AMDGPU also treats zero-sized
shared objects as dynamic LDS. Document this distinction in the CUDA guide
with a reference from the HIP documentation, and remove misleading wording
about struct overlays on dynamic shared memory.
Add CUDA and HIP CodeGen coverage for emitted types, shared address space,
and accesses at file and block scope. Extend semantic coverage to HIP and
check warning behavior in both non-RDC and RDC modes.
Validation: assertions-enabled Clang build; 125 SemaCUDA and shared-memory
CodeGen tests passed. Backend checks confirm sized PTX shared declarations
and 56 bytes of fixed LDS storage for the HIP test.
Pull Request: https://github.com/llvm/llvm-project/pull/183879
---
clang/docs/HIPSupport.md | 9 +++
.../clang/Basic/DiagnosticSemaKinds.td | 4 ++
clang/lib/Sema/SemaDeclAttr.cpp | 9 +--
clang/test/CodeGenCUDA/extern-shared.cu | 55 ++++++++++++++++
clang/test/SemaCUDA/extern-shared.cu | 65 +++++++++++++++----
llvm/docs/CompileCudaWithLLVM.md | 28 ++++++++
6 files changed, 155 insertions(+), 15 deletions(-)
create mode 100644 clang/test/CodeGenCUDA/extern-shared.cu
diff --git a/clang/docs/HIPSupport.md b/clang/docs/HIPSupport.md
index e2c9443145117f..0c961759f78f93 100644
--- a/clang/docs/HIPSupport.md
+++ b/clang/docs/HIPSupport.md
@@ -88,6 +88,15 @@ You can use `--offload-arch=native` to automatically detect the GPU architecture
clang++ --offload-arch=native -xhip sample.cpp -o sample
```
+## `extern __shared__` Variables
+
+HIP follows the same rules as CUDA for `extern __shared__` declarations.
+In non-RDC mode, nonzero-sized scalars, structs, and fixed-size arrays use fixed
+shared storage, while unsized arrays use dynamic shared memory. AMDGPU also
+treats zero-sized shared objects as dynamic LDS. See the
+[CUDA guide](https://llvm.org/docs/CompileCudaWithLLVM.html#extern-shared-variables)
+for examples and the optional `-Wcuda-extern-shared` warning.
+
## Path Setting for Dependencies
Compiling a HIP program depends on the HIP runtime and device library. The paths to the HIP runtime and device libraries
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index dcb5732c3c19e7..b4ed5dfebe125b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9852,6 +9852,10 @@ def err_cuda_vla : Error<
"cannot use variable-length arrays in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
def err_cuda_extern_shared : Error<"__shared__ variable %0 cannot be 'extern'">;
+def warn_cuda_extern_shared : Warning<
+ "'extern __shared__' variable %0 is not an incomplete array type; "
+ "NVCC accepts this but it is not standards-conforming">,
+ InGroup<DiagGroup<"cuda-extern-shared">>, DefaultIgnore;
def err_cuda_host_shared : Error<
"__shared__ local variables not allowed in "
"%select{__device__|__global__|__host__|__host__ __device__}0 functions">;
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 33b680c4d93d4c..a6307c3ef51cd2 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -5442,12 +5442,13 @@ static void handleConstantAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
static void handleSharedAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
const auto *VD = cast<VarDecl>(D);
- // extern __shared__ is only allowed on arrays with no length (e.g.
- // "int x[]").
+ // Accept extern __shared__ on complete types for NVCC compatibility, but
+ // optionally warn. In non-RDC mode nonzero-sized objects have fixed shared
+ // storage; incomplete arrays (e.g. "int x[]") use dynamic shared memory.
+ // AMDGPU also treats zero-sized objects as dynamic LDS.
if (!S.getLangOpts().GPURelocatableDeviceCode && VD->hasExternalStorage() &&
!isa<IncompleteArrayType>(VD->getType())) {
- S.Diag(AL.getLoc(), diag::err_cuda_extern_shared) << VD;
- return;
+ S.Diag(AL.getLoc(), diag::warn_cuda_extern_shared) << VD;
}
if (!S.CheckVarDeclSizeAddressSpace(VD, LangAS::cuda_shared))
return;
diff --git a/clang/test/CodeGenCUDA/extern-shared.cu b/clang/test/CodeGenCUDA/extern-shared.cu
new file mode 100644
index 00000000000000..7b9b0a45dcde0e
--- /dev/null
+++ b/clang/test/CodeGenCUDA/extern-shared.cu
@@ -0,0 +1,55 @@
+// NOTE: Do not autogenerate. These checks focus on shared types and accesses.
+// REQUIRES: nvptx-registered-target, amdgpu-registered-target
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -x hip -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
+
+// Non-RDC extern __shared__ declarations retain their declared types in shared
+// address space. Nonzero-sized complete types use fixed storage; unsized arrays
+// use dynamic shared memory. Check both file-scope and block-scope declarations.
+
+#include "Inputs/cuda.h"
+
+struct SharedData {
+ int x;
+ float y;
+};
+
+// CHECK-DAG: %struct.SharedData = type { i32, float }
+// CHECK-DAG: @scalar = external{{.*}} addrspace(3) global i32, align 4
+// CHECK-DAG: @data = external{{.*}} addrspace(3) global %struct.SharedData, align 4
+// CHECK-DAG: @fixed = external{{.*}} addrspace(3) global [4 x i32], align 4
+// CHECK-DAG: @dynamic = external{{.*}} addrspace(3) global [0 x i32], align 4
+// CHECK-DAG: @local_scalar = external{{.*}} addrspace(3) global i32, align 4
+// CHECK-DAG: @local_data = external{{.*}} addrspace(3) global %struct.SharedData, align 4
+// CHECK-DAG: @local_fixed = external{{.*}} addrspace(3) global [4 x i32], align 4
+// CHECK-DAG: @local_dynamic = external{{.*}} addrspace(3) global [0 x i32], align 4
+
+extern __shared__ int scalar;
+extern __shared__ SharedData data;
+extern __shared__ int fixed[4];
+extern __shared__ int dynamic[];
+
+// CHECK-LABEL: define{{.*}} void @_Z6kernelv(
+__global__ void kernel() {
+ // CHECK: store i32 1, ptr addrspacecast (ptr addrspace(3) @scalar to ptr)
+ scalar = 1;
+ // CHECK: store i32 2, ptr addrspacecast (ptr addrspace(3) @data to ptr)
+ data.x = 2;
+ // CHECK: store i32 3, ptr getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @fixed to ptr), i64 4)
+ fixed[1] = 3;
+ // CHECK: store i32 4, ptr getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @dynamic to ptr), i64 4)
+ dynamic[1] = 4;
+
+ extern __shared__ int local_scalar;
+ extern __shared__ SharedData local_data;
+ extern __shared__ int local_fixed[4];
+ extern __shared__ int local_dynamic[];
+ // CHECK: store i32 5, ptr addrspacecast (ptr addrspace(3) @local_scalar to ptr)
+ local_scalar = 5;
+ // CHECK: store i32 6, ptr addrspacecast (ptr addrspace(3) @local_data to ptr)
+ local_data.x = 6;
+ // CHECK: store i32 7, ptr getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @local_fixed to ptr), i64 4)
+ local_fixed[1] = 7;
+ // CHECK: store i32 8, ptr getelementptr inbounds nuw (i8, ptr addrspacecast (ptr addrspace(3) @local_dynamic to ptr), i64 4)
+ local_dynamic[1] = 8;
+}
diff --git a/clang/test/SemaCUDA/extern-shared.cu b/clang/test/SemaCUDA/extern-shared.cu
index bc4f10d3082c0d..5645573df1261d 100644
--- a/clang/test/SemaCUDA/extern-shared.cu
+++ b/clang/test/SemaCUDA/extern-shared.cu
@@ -1,31 +1,74 @@
// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -verify %s
// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s
-// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fgpu-rdc -verify=rdc %s
-// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -fcuda-is-device -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -Wcuda-extern-shared -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -fsyntax-only -Wundefined-internal -Wcuda-extern-shared -fcuda-is-device -fgpu-rdc -verify=rdc %s
+
+// RUN: %clang_cc1 -fsyntax-only -Wcuda-extern-shared -verify=warn %s
+// RUN: %clang_cc1 -fsyntax-only -Wcuda-extern-shared -fcuda-is-device -verify=warn %s
+
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wundefined-internal -verify %s
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wundefined-internal -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wundefined-internal -Wcuda-extern-shared -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wundefined-internal -Wcuda-extern-shared -fcuda-is-device -fgpu-rdc -verify=rdc %s
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wcuda-extern-shared -verify=warn %s
+// RUN: %clang_cc1 -x hip -fsyntax-only -Wcuda-extern-shared -fcuda-is-device -verify=warn %s
+
+// The warning is off by default in both CUDA and HIP.
+// expected-no-diagnostics
// Most of these declarations are fine in separate compilation mode.
#include "Inputs/cuda.h"
__device__ void foo() {
- extern __shared__ int x; // expected-error {{__shared__ variable 'x' cannot be 'extern'}}
+ // Scalar types.
+ extern __shared__ int x; // warn-warning {{'extern __shared__' variable 'x' is not an incomplete array type}}
+ extern __shared__ float f; // warn-warning {{'extern __shared__' variable 'f' is not an incomplete array type}}
+ extern __shared__ double d; // warn-warning {{'extern __shared__' variable 'd' is not an incomplete array type}}
+ extern __shared__ char c; // warn-warning {{'extern __shared__' variable 'c' is not an incomplete array type}}
+
+ // Incomplete array (the only standard-conforming form).
extern __shared__ int arr[]; // ok
- extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}}
- extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}}
- extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}}
+ extern __shared__ float farr[]; // ok
+ extern __shared__ double darr[]; // ok
+
+ // Fixed-size arrays and pointers are not incomplete arrays.
+#ifndef __HIP__
+ // HIP separately rejects zero-length arrays in device code.
+ extern __shared__ int arr0[0]; // warn-warning {{'extern __shared__' variable 'arr0' is not an incomplete array type}}
+#endif
+ extern __shared__ int arr1[1]; // warn-warning {{'extern __shared__' variable 'arr1' is not an incomplete array type}}
+ extern __shared__ float farr1[4]; // warn-warning {{'extern __shared__' variable 'farr1' is not an incomplete array type}}
+ extern __shared__ int* ptr; // warn-warning {{'extern __shared__' variable 'ptr' is not an incomplete array type}}
}
__host__ __device__ void bar() {
extern __shared__ int arr[]; // ok
- extern __shared__ int arr0[0]; // expected-error {{__shared__ variable 'arr0' cannot be 'extern'}}
- extern __shared__ int arr1[1]; // expected-error {{__shared__ variable 'arr1' cannot be 'extern'}}
- extern __shared__ int* ptr ; // expected-error {{__shared__ variable 'ptr' cannot be 'extern'}}
+#ifndef __HIP__
+ extern __shared__ int arr0[0]; // warn-warning {{'extern __shared__' variable 'arr0' is not an incomplete array type}}
+#endif
+ extern __shared__ int arr1[1]; // warn-warning {{'extern __shared__' variable 'arr1' is not an incomplete array type}}
+ extern __shared__ int* ptr; // warn-warning {{'extern __shared__' variable 'ptr' is not an incomplete array type}}
}
-extern __shared__ int global; // expected-error {{__shared__ variable 'global' cannot be 'extern'}}
+extern __shared__ int global; // warn-warning {{'extern __shared__' variable 'global' is not an incomplete array type}}
extern __shared__ int global_arr[]; // ok
-extern __shared__ int global_arr1[1]; // expected-error {{__shared__ variable 'global_arr1' cannot be 'extern'}}
+extern __shared__ int global_arr1[1]; // warn-warning {{'extern __shared__' variable 'global_arr1' is not an incomplete array type}}
+
+// A complete struct type uses fixed shared storage in non-RDC mode.
+struct ShmemData { int x; float y; };
+extern __shared__ ShmemData shmem; // warn-warning {{'extern __shared__' variable 'shmem' is not an incomplete array type}}
+__device__ void use_shmem() {
+ shmem.x = 1;
+ shmem.y = 2.0f;
+}
+
+// An incomplete array of structs uses dynamic shared memory.
+extern __shared__ ShmemData shmem_arr[]; // ok
+__device__ void use_shmem_arr() {
+ shmem_arr[0].x = 1;
+}
// Check that, iff we're not in rdc mode, extern __shared__ can appear in an
// anonymous namespace / in a static function without generating a warning
diff --git a/llvm/docs/CompileCudaWithLLVM.md b/llvm/docs/CompileCudaWithLLVM.md
index 577defef145ec4..c347cbeb1f2828 100644
--- a/llvm/docs/CompileCudaWithLLVM.md
+++ b/llvm/docs/CompileCudaWithLLVM.md
@@ -212,6 +212,34 @@ dialects of the language. Below, we describe some of the differences.
This section is painful; hopefully you can skip this section and live your life
blissfully unaware.
+### `extern __shared__` Variables
+
+Clang accepts `extern __shared__` declarations of scalars, structs, and
+fixed-size arrays for compatibility with nvcc. In non-RDC mode
+(`-fno-gpu-rdc`), nonzero-sized objects use fixed shared storage sized according
+to their declared types. Unsized arrays use dynamic shared memory,
+whose size is specified at kernel launch. AMDGPU also treats zero-sized
+shared objects as dynamic LDS. A nonzero-sized complete struct declared
+`extern __shared__` does not overlay dynamic shared memory.
+
+```cuda
+struct SharedData { int x; float y; };
+
+__global__ void kernel() {
+ extern __shared__ int scalar; // Fixed shared storage: one int.
+ extern __shared__ SharedData data; // Fixed shared storage: one SharedData.
+ extern __shared__ int fixed[16]; // Fixed shared storage: 16 ints.
+ extern __shared__ int dynamic[]; // Dynamic shared memory.
+}
+```
+
+`-Wcuda-extern-shared` enables a warning for non-RDC `extern __shared__`
+declarations whose types are not incomplete arrays. This warning is disabled
+by default. In RDC mode (`-fgpu-rdc`), these declarations are accepted without
+this warning and follow the usual rules for external device variables.
+
+These rules also apply to HIP. See [HIP Support](https://clang.llvm.org/docs/HIPSupport.html).
+
### Compilation Models
Most of the differences between clang and nvcc stem from the different
More information about the cfe-commits
mailing list