[clang] [OpenCL] Add clang internal extension __cl_clang_non_kernel_scope_local_memory (PR #176726)
Wenju He via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 19 02:45:55 PST 2026
https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/176726
>From da646fb0e0058aeb9e23dc9ead944cb42a8f9da3 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Mon, 19 Jan 2026 11:35:48 +0100
Subject: [PATCH 1/2] [OpenCL] Add clang internal extension
__cl_clang_non_kernel_scope_local_memory
OpenCL spec restricts that variable in local address space can only be
declared at kernel function scope.
Aad a Clang internal extension __cl_clang_non_kernel_scope_local_memory
to lift the restriction.
With this relaxation, targets can force-inline non-kernel functions that
declare local memory - so static local allocations are visible at kernel
scope - or pass a kernel-allocated local buffer to those functions via
an implicit argument.
Motivation: support local memory allocation in libclc's implementation
of work-group collective built-ins, see example at:
https://github.com/intel/llvm/blob/41455e305117/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives_helpers.ll
https://github.com/intel/llvm/blob/41455e305117/libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl#L182
---
clang/docs/LanguageExtensions.rst | 44 +++++++++++++++++++
.../include/clang/Basic/OpenCLExtensions.def | 1 +
clang/lib/Sema/SemaDecl.cpp | 13 +++++-
.../CodeGenOpenCL/local-non-kernel-scope.cl | 19 ++++++++
clang/test/SemaOpenCL/extension-version.cl | 5 +++
clang/test/SemaOpenCL/storageclass.cl | 31 +++++++++----
6 files changed, 102 insertions(+), 11 deletions(-)
create mode 100644 clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index 228f7bf89ddde..a3498576be725 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -2840,6 +2840,50 @@ between the host and device is known to be compatible.
);
#pragma OPENCL EXTENSION __cl_clang_non_portable_kernel_param_types : disable
+``__cl_clang_non_kernel_scope_local_memory``
+----------------------------------------------
+
+This extension allows declaring variables in the local address space within
+non-kernel functions or nested scopes within a kernel, using regular OpenCL
+extension pragma mechanism detailed in `the OpenCL Extension Specification,
+section 1.2
+<https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_Ext.html#extensions-overview>`_.
+
+This relaxes the `Declaration Scopes and Variable Types
+<https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_C.html#_usage_for_declaration_scopes_and_variable_types>`_
+rule that limits local-address-space variable declarations to the outermost scope
+of a kernel function only.
+
+With this relaxation, targets can force-inline non-kernel functions that declare
+local memory - so static local allocations are visible at kernel scope - or pass
+a kernel-allocated local buffer to those functions via an implicit argument.
+
+.. code-block:: c++
+
+ #pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+ kernel void kernel1(...)
+ {
+ {
+ local float a; // compiled - no diagnostic generated
+ }
+ }
+ void foo()
+ {
+ local float c; // compiled - no diagnostic generated
+ }
+
+ #pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : disable
+ kernel void kernel2(...)
+ {
+ {
+ local float a; // error - non-kernel function variable cannot be declared in local address space.
+ }
+ }
+ void bar()
+ {
+ local float c; // error - non-kernel function variable cannot be declared in local address space.
+ }
+
Remove address space builtin function
-------------------------------------
diff --git a/clang/include/clang/Basic/OpenCLExtensions.def b/clang/include/clang/Basic/OpenCLExtensions.def
index d6c0b585d1809..3ce81e6a769d7 100644
--- a/clang/include/clang/Basic/OpenCLExtensions.def
+++ b/clang/include/clang/Basic/OpenCLExtensions.def
@@ -131,6 +131,7 @@ OPENCL_GENERIC_EXTENSION(__opencl_c_work_group_collective_functions, false, 200,
OPENCL_EXTENSION(cl_clang_storage_class_specifiers, true, 100)
OPENCL_EXTENSION(__cl_clang_function_pointers, true, 100)
OPENCL_EXTENSION(__cl_clang_variadic_functions, true, 100)
+OPENCL_EXTENSION(__cl_clang_non_kernel_scope_local_memory, true, 100)
OPENCL_EXTENSION(__cl_clang_non_portable_kernel_param_types, true, 100)
OPENCL_EXTENSION(__cl_clang_bitfields, true, 100)
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index ae779d6830d9b..8e602a7c6d4b9 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8948,8 +8948,17 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
NewVD->setInvalidDecl();
return;
}
- if (T.getAddressSpace() == LangAS::opencl_constant ||
- T.getAddressSpace() == LangAS::opencl_local) {
+ // When this extension is enabled, 'local' variables are permitted in
+ // non-kernel functions and within nested scopes of kernel functions,
+ // bypassing standard OpenCL address space restrictions.
+ bool AllowNonKernelLocal =
+ T.getAddressSpace() == LangAS::opencl_local &&
+ getOpenCLOptions().isAvailableOption(
+ "__cl_clang_non_kernel_scope_local_memory", getLangOpts());
+ if (AllowNonKernelLocal) {
+ // Direct pass: No further diagnostics needed for this specific case.
+ } else if (T.getAddressSpace() == LangAS::opencl_constant ||
+ T.getAddressSpace() == LangAS::opencl_local) {
FunctionDecl *FD = getCurFunctionDecl();
// OpenCL v1.1 s6.5.2 and s6.5.3: no local or constant variables
// in functions.
diff --git a/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl b/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
new file mode 100644
index 0000000000000..0bdec8fd7fdd7
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 %s -triple spir64 -disable-llvm-passes -emit-llvm -o - | FileCheck %s
+
+#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+
+void func(local int*);
+
+void bar() {
+ // CHECK: @bar.i = internal addrspace(3) global i32 undef, align 4
+ local int i;
+ func(&i);
+}
+
+__kernel void foo(void) {
+ // CHECK: @foo.i = internal addrspace(3) global i32 undef, align 4
+ {
+ local int i;
+ func(&i);
+ }
+}
diff --git a/clang/test/SemaOpenCL/extension-version.cl b/clang/test/SemaOpenCL/extension-version.cl
index b24c1b4bb6272..3b81769970fbd 100644
--- a/clang/test/SemaOpenCL/extension-version.cl
+++ b/clang/test/SemaOpenCL/extension-version.cl
@@ -27,6 +27,11 @@
#endif
#pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable
+#ifndef __cl_clang_non_kernel_scope_local_memory
+#error "Missing __cl_clang_non_kernel_scope_local_memory define"
+#endif
+#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+
#ifndef cl_khr_fp16
#error "Missing cl_khr_fp16 define"
#endif
diff --git a/clang/test/SemaOpenCL/storageclass.cl b/clang/test/SemaOpenCL/storageclass.cl
index 4b9d6e9dd4f2d..d500391e2b4ce 100644
--- a/clang/test/SemaOpenCL/storageclass.cl
+++ b/clang/test/SemaOpenCL/storageclass.cl
@@ -1,12 +1,12 @@
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_generic_address_space
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__opencl_c_generic_address_space
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_generic_address_space
-// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__opencl_c_generic_address_space
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2 -cl-ext=+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__opencl_c_generic_address_space,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_non_kernel_scope_local_memory
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__opencl_c_generic_address_space,+__cl_clang_non_kernel_scope_local_memory
static constant int G1 = 0;
constant int G2 = 0;
@@ -278,3 +278,16 @@ void f(void) {
#endif
#endif
}
+
+void f_local(void) {
+#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+ local int L2;
+ {
+ local int L2;
+ }
+#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : disable
+ local int L2; // expected-error{{non-kernel function variable cannot be declared in local address space}}
+ {
+ local int L2; // expected-error{{non-kernel function variable cannot be declared in local address space}}
+ }
+}
>From 6f012f3c840180eec958515a52490c941d99de08 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Mon, 19 Jan 2026 11:45:43 +0100
Subject: [PATCH 2/2] update per coplit review comment
---
clang/docs/LanguageExtensions.rst | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index a3498576be725..c835092129a80 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -2876,12 +2876,12 @@ a kernel-allocated local buffer to those functions via an implicit argument.
kernel void kernel2(...)
{
{
- local float a; // error - non-kernel function variable cannot be declared in local address space.
+ local float a; // error - variables in the local address space can only be declared in the outermost scope of a kernel function
}
}
void bar()
{
- local float c; // error - non-kernel function variable cannot be declared in local address space.
+ local float c; // error - non-kernel function variable cannot be declared in local address space
}
Remove address space builtin function
More information about the cfe-commits
mailing list