[clang] [OpenCL] Add clang internal extension __cl_clang_non_kernel_scope_local_memory (PR #176726)
via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 19 02:41:26 PST 2026
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Wenju He (wenju-he)
<details>
<summary>Changes</summary>
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
---
Full diff: https://github.com/llvm/llvm-project/pull/176726.diff
6 Files Affected:
- (modified) clang/docs/LanguageExtensions.rst (+44)
- (modified) clang/include/clang/Basic/OpenCLExtensions.def (+1)
- (modified) clang/lib/Sema/SemaDecl.cpp (+11-2)
- (added) clang/test/CodeGenOpenCL/local-non-kernel-scope.cl (+19)
- (modified) clang/test/SemaOpenCL/extension-version.cl (+5)
- (modified) clang/test/SemaOpenCL/storageclass.cl (+22-9)
``````````diff
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}}
+ }
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/176726
More information about the cfe-commits
mailing list