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 <[email protected]> Date: Mon, 19 Jan 2026 11:35:48 +0100 Subject: [PATCH 1/3] [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 <[email protected]> Date: Mon, 19 Jan 2026 11:45:43 +0100 Subject: [PATCH 2/3] 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 <[email protected]> Date: Tue, 20 Jan 2026 03:52:06 +0100 Subject: [PATCH 3/3] 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}} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
