[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 19:13:57 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/4] [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/4] 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
>From aaa41fb5e834109058fd808f1cce84ad771b0ea8 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 20 Jan 2026 03:52:06 +0100
Subject: [PATCH 3/4] rename to __cl_clang_local_memory_all_scopes
---
clang/docs/LanguageExtensions.rst | 22 +++++++++----------
.../include/clang/Basic/OpenCLExtensions.def | 2 +-
clang/lib/Sema/SemaDecl.cpp | 2 +-
.../CodeGenOpenCL/local-non-kernel-scope.cl | 2 +-
clang/test/SemaOpenCL/extension-version.cl | 6 ++---
clang/test/SemaOpenCL/storageclass.cl | 22 +++++++++----------
6 files changed, 28 insertions(+), 28 deletions(-)
diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index c835092129a80..b878a6c298ace 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -2840,27 +2840,27 @@ 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``
+``__cl_clang_local_memory_all_scopes``
----------------------------------------------
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
+any 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 scope
-of a kernel function only.
+rule that limits local-address-space variable declarations to the outermost
+compound statement inside the body of the kernel function.
-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.
+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_non_kernel_scope_local_memory : enable
+ #pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : enable
kernel void kernel1(...)
{
{
@@ -2872,7 +2872,7 @@ a kernel-allocated local buffer to those functions via an implicit argument.
local float c; // compiled - no diagnostic generated
}
- #pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : disable
+ #pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : disable
kernel void kernel2(...)
{
{
diff --git a/clang/include/clang/Basic/OpenCLExtensions.def b/clang/include/clang/Basic/OpenCLExtensions.def
index 3ce81e6a769d7..270091029ad5f 100644
--- a/clang/include/clang/Basic/OpenCLExtensions.def
+++ b/clang/include/clang/Basic/OpenCLExtensions.def
@@ -131,7 +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_local_memory_all_scopes, 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 8e602a7c6d4b9..9755da9897be8 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8954,7 +8954,7 @@ void Sema::CheckVariableDeclarationType(VarDecl *NewVD) {
bool AllowNonKernelLocal =
T.getAddressSpace() == LangAS::opencl_local &&
getOpenCLOptions().isAvailableOption(
- "__cl_clang_non_kernel_scope_local_memory", getLangOpts());
+ "__cl_clang_local_memory_all_scopes", getLangOpts());
if (AllowNonKernelLocal) {
// Direct pass: No further diagnostics needed for this specific case.
} else if (T.getAddressSpace() == LangAS::opencl_constant ||
diff --git a/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl b/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
index 0bdec8fd7fdd7..50193896fccdf 100644
--- a/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
+++ b/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
@@ -1,6 +1,6 @@
// 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
+#pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : enable
void func(local int*);
diff --git a/clang/test/SemaOpenCL/extension-version.cl b/clang/test/SemaOpenCL/extension-version.cl
index 3b81769970fbd..f99b451eba301 100644
--- a/clang/test/SemaOpenCL/extension-version.cl
+++ b/clang/test/SemaOpenCL/extension-version.cl
@@ -27,10 +27,10 @@
#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"
+#ifndef __cl_clang_local_memory_all_scopes
+#error "Missing __cl_clang_local_memory_all_scopes define"
#endif
-#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+#pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : enable
#ifndef cl_khr_fp16
#error "Missing cl_khr_fp16 define"
diff --git a/clang/test/SemaOpenCL/storageclass.cl b/clang/test/SemaOpenCL/storageclass.cl
index d500391e2b4ce..67afe4c042ed3 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 -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
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL1.2 -cl-ext=+__cl_clang_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__cl_clang_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL3.0 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_local_memory_all_scopes
+// 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_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__cl_clang_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_program_scope_global_variables,+__cl_clang_local_memory_all_scopes
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=clc++2021 -cl-ext=-all,+__opencl_c_generic_address_space,+__cl_clang_local_memory_all_scopes
+// 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_local_memory_all_scopes
static constant int G1 = 0;
constant int G2 = 0;
@@ -280,12 +280,12 @@ void f(void) {
}
void f_local(void) {
-#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : enable
+#pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : enable
local int L2;
{
local int L2;
}
-#pragma OPENCL EXTENSION __cl_clang_non_kernel_scope_local_memory : disable
+#pragma OPENCL EXTENSION __cl_clang_local_memory_all_scopes : 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 d9e71fcec98ca13b01b6f438d9a0ec60ae3db6f1 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 20 Jan 2026 04:13:40 +0100
Subject: [PATCH 4/4] rename test to local-scope.cl
---
.../CodeGenOpenCL/{local-non-kernel-scope.cl => local-scope.cl} | 0
1 file changed, 0 insertions(+), 0 deletions(-)
rename clang/test/CodeGenOpenCL/{local-non-kernel-scope.cl => local-scope.cl} (100%)
diff --git a/clang/test/CodeGenOpenCL/local-non-kernel-scope.cl b/clang/test/CodeGenOpenCL/local-scope.cl
similarity index 100%
rename from clang/test/CodeGenOpenCL/local-non-kernel-scope.cl
rename to clang/test/CodeGenOpenCL/local-scope.cl
More information about the cfe-commits
mailing list