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}} + } +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
