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