[clang] c03d0fe - [OpenCL] Add clang internal extension __cl_clang_function_scope_local_variables (#176726)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Jan 25 16:13:26 PST 2026
Author: Wenju He
Date: 2026-01-26T08:13:22+08:00
New Revision: c03d0fe672b2506d22d101cb7704996f4ffda3e4
URL: https://github.com/llvm/llvm-project/commit/c03d0fe672b2506d22d101cb7704996f4ffda3e4
DIFF: https://github.com/llvm/llvm-project/commit/c03d0fe672b2506d22d101cb7704996f4ffda3e4.diff
LOG: [OpenCL] Add clang internal extension __cl_clang_function_scope_local_variables (#176726)
OpenCL spec restricts that variable in local address space can only be
declared at kernel function scope.
Add a Clang internal extension __cl_clang_function_scope_local_variables
to lift the restriction.
To expose static local allocations at kernel scope, targets can either
force-inline non-kernel functions that declare local memory 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
Right now this is a Clang-only OpenCL extension intended for compiling
OpenCL libraries with Clang. It could be proposed as a standard OpenCL
extension in the future.
Added:
clang/test/CodeGenOpenCL/local-scope.cl
Modified:
clang/docs/LanguageExtensions.rst
clang/include/clang/Basic/OpenCLExtensions.def
clang/lib/Sema/SemaDecl.cpp
clang/test/SemaOpenCL/extension-version.cl
clang/test/SemaOpenCL/storageclass.cl
Removed:
################################################################################
diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index dc0f838e874a0..0adfaebf24581 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -2884,6 +2884,50 @@ between the host and device is known to be compatible.
);
#pragma OPENCL EXTENSION __cl_clang_non_portable_kernel_param_types : disable
+``__cl_clang_function_scope_local_variables``
+----------------------------------------------
+
+This extension allows declaring variables in the local address space within
+function scope, including 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
+compound statement inside the body of the kernel function.
+
+To expose static local allocations at kernel scope, targets can either force-
+inline non-kernel functions that declare local memory or pass a kernel-allocated
+local buffer to those functions via an implicit argument.
+
+.. code-block:: c++
+
+ #pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : 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_function_scope_local_variables : disable
+ kernel void kernel2(...)
+ {
+ {
+ 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
+ }
+
Remove address space builtin function
-------------------------------------
diff --git a/clang/include/clang/Basic/OpenCLExtensions.def b/clang/include/clang/Basic/OpenCLExtensions.def
index d6c0b585d1809..8113cdb44f367 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_function_scope_local_variables, 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 066acc3424c8f..3b14abe993034 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8975,8 +8975,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 AllowFunctionScopeLocalVariables =
+ T.getAddressSpace() == LangAS::opencl_local &&
+ getOpenCLOptions().isAvailableOption(
+ "__cl_clang_function_scope_local_variables", getLangOpts());
+ if (AllowFunctionScopeLocalVariables) {
+ // 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-scope.cl b/clang/test/CodeGenOpenCL/local-scope.cl
new file mode 100644
index 0000000000000..7b61a32e95c96
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/local-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_function_scope_local_variables : 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..1dd3839a5c603 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_function_scope_local_variables
+#error "Missing __cl_clang_function_scope_local_variables define"
+#endif
+#pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : 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..fc6518b2d42df 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_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__cl_clang_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_function_scope_local_variables
+// 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_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__cl_clang_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_function_scope_local_variables
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_function_scope_local_variables
+// 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_function_scope_local_variables
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_function_scope_local_variables : enable
+ local int L2;
+ {
+ local int L2;
+ }
+#pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : 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}}
+ }
+}
More information about the cfe-commits
mailing list