https://github.com/yuxuanchen1997 updated https://github.com/llvm/llvm-project/pull/183879
>From 7b578758cb889cd8863085cb06fe43f1b536a0d6 Mon Sep 17 00:00:00 2001 From: Yuxuan Chen <[email protected]> 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, complete types use fixed shared storage; only unsized arrays use dynamic shared memory. 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 | 8 +++ .../clang/Basic/DiagnosticSemaKinds.td | 4 ++ clang/lib/Sema/SemaDeclAttr.cpp | 8 +-- clang/test/CodeGenCUDA/extern-shared.cu | 54 +++++++++++++++ clang/test/SemaCUDA/extern-shared.cu | 65 +++++++++++++++---- llvm/docs/CompileCudaWithLLVM.md | 27 ++++++++ 6 files changed, 151 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..68d127b597f0bb 100644 --- a/clang/docs/HIPSupport.md +++ b/clang/docs/HIPSupport.md @@ -88,6 +88,14 @@ 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, scalars, structs, and fixed-size arrays use fixed shared +storage, while unsized arrays use dynamic shared memory. 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..af3ca184ffd6fa 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5442,12 +5442,12 @@ 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 these have fixed shared storage; only + // incomplete arrays (e.g. "int x[]") use dynamic shared memory. 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..3153258d229d89 --- /dev/null +++ b/clang/test/CodeGenCUDA/extern-shared.cu @@ -0,0 +1,54 @@ +// 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 amdgcn-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. Complete types use fixed storage; only 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..0240645ca43119 100644 --- a/llvm/docs/CompileCudaWithLLVM.md +++ b/llvm/docs/CompileCudaWithLLVM.md @@ -212,6 +212,33 @@ 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`), these declarations use fixed shared storage sized according +to their declared types. Only unsized arrays use dynamic shared memory, +whose size is specified at kernel launch. A 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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
