https://github.com/tahonermann created https://github.com/llvm/llvm-project/pull/170602
Clang's emerging support for the SYCL 2020 specification includes new driver command line options, predefined macros, and several C++11-style attributes that enable core language features intended to support implementation of SYCL run-time libraries. Included is documentation intended to provide an architectural overview with a target audience of SYCL implementers. This update discards previous documentation from an earlier SYCL implementation effort that is unlikely to be correct or relevant for the new SYCL implementation effort. >From 790b9374349e1621bf6514e61834084b750540b6 Mon Sep 17 00:00:00 2001 From: Tom Honermann <[email protected]> Date: Thu, 20 Nov 2025 08:55:37 -0800 Subject: [PATCH] [SYCL] Documentation for Clang driver and core language support. Clang's emerging support for the SYCL 2020 specification includes new driver command line options, predefined macros, and several C++11-style attributes that enable core language features intended to support implementation of SYCL run-time libraries. Included is documentation intended to provide an architectural overview with a target audience of SYCL implementers. This update discards previous documentation from an earlier SYCL implementation effort that is unlikely to be correct or relevant for the new SYCL implementation effort. --- clang/docs/SYCLSupport.rst | 615 +++++++++++++++++++++++++++++++------ 1 file changed, 519 insertions(+), 96 deletions(-) diff --git a/clang/docs/SYCLSupport.rst b/clang/docs/SYCLSupport.rst index 364d95a31047e..d0ce5ab885f1a 100644 --- a/clang/docs/SYCLSupport.rst +++ b/clang/docs/SYCLSupport.rst @@ -1,119 +1,542 @@ -============================================= -SYCL Compiler and Runtime architecture design -============================================= +============ +SYCL Support +============ .. contents:: :local: Introduction ============ +The `SYCL 2020 specification <SYCL-2020_>`_ defines a single-source programming +model and C++ run-time library interface for writing portable programs that +support heterogeneous devices including GPUs, CPUs, and accelerators. +The specification is intended to allow for a wide range of implementation +possibilities, examples of which include: + +- A SYCL run-time library written in standard C++ that executes kernels on a + homogeneous set of host and device processors, each of which can execute + common compiled code from shared memory. +- A SYCL run-time library that executes kernels on a heterogeneous set of + device processors for which each kernel is pre-compiled for each supported + device processor (Ahead-Of-Time (AOT) compilation) or for a family of device + processors (Just-In-Time (JIT) compilation). + +Since Clang is a conforming implementation of the C++ standard, no additional +features are required for support of the first implementation strategy. +This document details the core language features Clang provides for use by +SYCL run-time libraries that use the second implementation strategy. + +.. _SYCL-2020: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html + + +Example Usage +============= +SYCL is designed as an extension of C++ rather than as a distinct programming +language. +SYCL support is enabled with the `-fsycl <opt-fsycl_>`_ option. + +.. code-block:: sh + + clang++ -c -fsycl source-file.cpp + +The choice of which target devices will be supported is made at compile time. +By default, SYCL source files will be compiled with support for a host target +dependent set of target devices. +For example, when compiling for a ``x86_64-unknown-linux-gnu`` host target, +target support will be enabled for ``spirv64-unknown-unknown`` devices. +The set of supported target devices can be specified via a comma separated list +of target triples with the `--offload-targets= <opt-offload-targets_>`_ option. +The following Clang invocation enables support for AMD, NVIDIA, and Intel GPU +targets. + +.. code-block:: sh + + clang++ -c -fsycl \ + --offload-targets=amdgcn-amd-amdhsa,nvptx64-nvidia-cuda,spirv64-unknown-unknown \ + source-file.cpp + +Object files built with the `-fsycl <opt-fsycl_>`_ option contain device +images that require additional processing at link time. +Programs linked with such object files must also be linked using the +``clang++`` driver and the `-fsycl <opt-fsycl_>`_ option. + +.. code-block:: sh + + clang++ -fsycl example.o source-file.o -o example + +.. _opt-fsycl: + https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-fsycl +.. _opt-offload-targets: + https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-offload-targets + + +Compilation Model +================= +`SYCL 2020 section 5.1, "Offline compilation of SYCL source files" <SYCL-2020-5.1_>`_ +acknowledges two compilation models. + +- Single-source Multiple Compiler Pass (`SMCP`_) describes a compilation model + in which source code is separately parsed and analyzed for the host target + and each device target. + +- Single-source Single Compiler Pass (`SSCP`_) describes a compilation model + in which source code is parsed and analyzed once with code generation + performed separately for the host target and each device target. + +Clang only supports the `SMCP`_ compilation model currently, but the SYCL +language support features have been designed to allow for support of the +`SSCP`_ compilation model to be added in the future. + +By default, SYCL source files are compiled for the host target and for each +device target. +In some cases, it is useful to restrict compilation to just the host target or +just the device targets; the `-fsycl-host-only <opt-fsycl-host-only_>`_ and +`-fsycl-device-only <opt-fsycl-device-only_>`_ options are available for these +purposes. + +.. _SMCP: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:smcp +.. _SSCP: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:sscp +.. _SYCL-2020-5.1: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_offline_compilation_of_sycl_source_files +.. _opt-fsycl-host-only: + https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-offload-host-only +.. _opt-fsycl-device-only: + https://clang.llvm.org/docs/ClangCommandLineReference.html#cmdoption-clang-offload-device-only + + +Supported Targets +================= +Support for SYCL is still in the implementation phase, but all targets +supported by the `--offload-targets= <opt-offload-targets_>`_ option +are intended to eventually be supported. -This document describes the architecture of the SYCL compiler and runtime -library. More details are provided in -`external document <https://github.com/intel/llvm/blob/sycl/sycl/doc/design/CompilerAndRuntimeDesign.md>`_\ , -which are going to be added to clang documentation in the future. - -Address space handling -====================== - -The SYCL specification represents pointers to disjoint memory regions using C++ -wrapper classes on an accelerator to enable compilation with a standard C++ -toolchain and a SYCL compiler toolchain. Section 3.8.2 of SYCL 2020 -specification defines -`memory model <https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_sycl_device_memory_model>`_\ , -section 4.7.7 - `address space classes <https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_address_space_classes>`_ -and section 5.9 covers `address space deduction <https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#_address_space_deduction>`_. -The SYCL specification allows two modes of address space deduction: "generic as -default address space" (see section 5.9.3) and "inferred address space" (see -section 5.9.4). Current implementation supports only "generic as default address -space" mode. - -SYCL borrows its memory model from OpenCL however SYCL doesn't perform -the address space qualifier inference as detailed in -`OpenCL C v3.0 6.7.8 <https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#addr-spaces-inference>`_. - -The default address space is "generic-memory", which is a virtual address space -that overlaps the global, local, and private address spaces. SYCL mode enables -following conversions: - -- explicit conversions to/from the default address space from/to the address - space-attributed type -- implicit conversions from the address space-attributed type to the default - address space -- explicit conversions to/from the global address space from/to the - ``__attribute__((opencl_global_device))`` or - ``__attribute__((opencl_global_host))`` address space-attributed type -- implicit conversions from the ``__attribute__((opencl_global_device))`` or - ``__attribute__((opencl_global_host))`` address space-attributed type to the - global address space - -All named address spaces are disjoint and sub-sets of default address space. - -The SPIR target allocates SYCL namespace scope variables in the global address -space. - -Pointers to default address space should get lowered into a pointer to a generic -address space (or flat to reuse more general terminology). But depending on the -allocation context, the default address space of a non-pointer type is assigned -to a specific address space. This is described in -`common address space deduction rules <https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:commonAddressSpace>`_ -section. - -This is also in line with the behaviour of CUDA (`small example -<https://godbolt.org/z/veqTfo9PK>`_). - -``multi_ptr`` class implementation example: + +Predefined Macros +================= +`SYCL 2020 section 5.6, "Preprocessor directives and macros" <SYCL-2020-5.6_>`_ +specifies macros that a SYCL implementation is required to provide. +Most such macros are defined by the SYCL run-time library and require inclusion +of the ``<sycl/sycl.hpp>`` header file. +The following macros are conditionally predefined by the compiler. + +.. list-table:: + :header-rows: 1 + + * - Macro + - Description + * - ``__SYCL_DEVICE_ONLY__`` + - Predefined by a `SMCP`_ implementation during device compilation (but not + during host compilation). + * - ``__SYCL_SINGLE_SOURCE__`` + - Predefined by a `SSCP`_ implementation during (host and device) + compilation. + +Since Clang only supports the `SMCP`_ compilation model currently, the +``__SYCL_SINGLE_SOURCE__`` macro is never predefined. + +.. _SYCL-2020-5.6: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_preprocessor_directives_and_macros + + +Language Support +================ +`SYCL 2020 section 3.12.3, "Library-only implementation" <SYCL-2020-3.12.13_>`_ +notes the intent that the SYCL specification be implementable as a C++ library +with no requirements beyond a compiler that conforms to the C++17 standard. +The SYCL specification therefore does not specify extensions to the C++ core +language and a library-only implementation will work with Clang without any +core language extensions. +Clang provides the features described in this section to facilitate capabilities +that are not possible with a library-only SYCL implementation. + +.. _SYCL-2020-3.12.13: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_library_only_implementation + + +.. _sect-sycl_kernel_entry_point: + +The ``[[clang::sycl_kernel_entry_point]]`` Attribute +---------------------------------------------------- +This attribute is intended for use in the implementation of SYCL run-time +libraries and should not be used directly by programmers. + +The `sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ attribute +coordinates interaction between Clang and the SYCL run-time library to +facilitate code generation and the execution of a SYCL kernel on a device +that potentially uses an instruction set architecture different from the host. +Consider the following call to the ``sycl::handler::single_task()`` SYCL +kernel invocation function. .. code-block:: C++ - // check that SYCL mode is ON and we can use non-standard decorations - #if defined(__SYCL_DEVICE_ONLY__) - // GPU/accelerator implementation - template <typename T, address_space AS> class multi_ptr { - // DecoratedType applies corresponding address space attribute to the type T - // DecoratedType<T, global_space>::type == "__attribute__((opencl_global)) T" - // See sycl/include/CL/sycl/access/access.hpp for more details - using pointer_t = typename DecoratedType<T, AS>::type *; + struct KN; + void f(sycl::handler &h, sycl::stream &sout, int i) { + h.single_task<KN>([=] { + sout << "The value of i is " << i << "\n"; + }); + } + +The SYCL kernel is defined by the lambda expression passed to the +``single_task()`` function and is identified by the ``KN`` type passed as the +first template argument. +See +`SYCL 2020 section 4.9.4.2, "SYCL functions for invoking kernels" <SYCL-2020-4.9.4.2_>`_ +and +`SYCL 2020 section 5.2, "Naming of kernels" <SYCL-2020-5.2_>`_ +for more details. + +The `SMCP`_ and `SSCP`_ compilation models require that code generation be +performed for each SYCL kernel for each target device. +In order for Clang to perform that code generation, it needs to be informed +that a SYCL kernel invocation is present. +The `sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ attribute +provides the means for the SYCL run-time library, which provides the +definition of all SYCL kernel invocation functions, to inform Clang of a SYCL +kernel invocation. +This is accomplished by, in the definition of a SYCL kernel invocation function, +including a call to a function declared with the attribute. +For example: + +.. code-block:: C++ - pointer_t m_Pointer; + namespace sycl { + class handler { + template <typename KernelName, typename KernelType> + [[clang::sycl_kernel_entry_point(KernelName)]] + void kernel_entry_point(KernelType kernelFunc) { + kernelFunc(); + } public: - pointer_t get() { return m_Pointer; } - T& operator* () { return *reinterpret_cast<T*>(m_Pointer); } + template <typename KernelName, typename KernelType> + void single_task(const KernelType &kernelFunc) { + kernel_entry_point<KernelName>(kernelFunc); + } + }; } - #else - // CPU/host implementation - template <typename T, address_space AS> class multi_ptr { - T *m_Pointer; // regular undecorated pointer - public: - T *get() { return m_Pointer; } - T& operator* () { return *m_Pointer; } + +The arguments of the call to ``kernel_entry_point()`` constitute the parameters +of a SYCL kernel. +The body of the ``kernel_entry_point()`` function contains the statements +required to execute the SYCL kernel (e.g., an invocation of the call operator of +the SYCL kernel object as in this example). +The call to ``kernel_entry_point()`` in ``single_task()`` establishes a common +point of SYCL kernel invocation for Clang and the SYCL run-time library. +With that point established, the tasks required to actually execute a SYCL +kernel are delegated according to the following division of responsibilities. + +Clang is responsible for: + +- Validating that all kernel argument types (e.g., the deduced parameter types + of the ``kernel_entry_point()`` function above) satisfy the requirements + specified in + `SYCL 2020 section 4.12.4, "Rules for parameter passing to kernels" <SYCL-2020-4.12.4_>`_. +- Informing the SYCL run-time library of the presence of subobjects of SYCL + types that require special handling within kernel arguments. +- Generating an offload kernel entry point function for each SYCL kernel for + each target device, generating a name for it derived from the SYCL kernel + name, and informing the SYCL run-time library of the generated name. + +The SYCL run-time library is responsible for: + +- Selecting a device on which to execute the kernel. +- Copying the SYCL kernel object and any other kernel arguments to the device. +- Informing Clang of additional parameters required for the offload kernel + entry point based on the presence of subobjects of SYCL types that require + special handling within kernel arguments. +- Scheduling execution of the offload kernel entry point function on the + selected device. + +The SYCL run-time library tasks are expected to be performed in conjunction +with an offload backend such as liboffload, OpenCL, CUDA, Hip, or Level Zero; +their details are out of scope for this document. + +The above division of responsibilities requires coordination. +The call to a function declared with the +`sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ attribute causes +two primary side effects: + +- The generation of an offload kernel entry point function. +- An implicit call to a SYCL run-time library provided template named + ``sycl_kernel_launch`` (which may be a function template or a variable + template of a type with a member call operator). + +The offload kernel entry point function is generated with a target dependent +calling convention for each device target. +The function parameters and function body are initially copied from the function +declared with the `sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ +attribute, but may be augmented by information provided by the SYCL run-time +library as described below. +The function name is an implementation detail subject to change, but +incorporates the SYCL kernel name in order to ensure that a unique name is +deterministically generated for each SYCL kernel. + +The call to the ``sycl_kernel_launch`` template effectively replaces the call +to the `sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ attributed +function. +This implicit call serves several purposes: + +- It informs the SYCL run-time library of the name of the offload kernel entry + point function to be used to execute the kernel on the selected device. +- It informs the SYCL run-time library of the presence of subobjects of the + kernel arguments that require special handling. + +See the +:ref:`sycl_special_kernel_parameter <sect-sycl_special_kernel_parameter>` +section regarding SYCL types that require special handling in kernel arguments. + +The call to the ``sycl_kernel_launch`` template passes the name of the generated +offload kernel entry point function, the kernel arguments, and, for each +subobject of a kernel argument that requires special handling, a reference to +that subobject. +For reasons explained further below, the ``sycl_kernel_launch`` template needs +to know which of its arguments correspond to direct kernel arguments and which +correspond to references to special subobjects within the direct kernel +arguments. +Because there may be multiple kernel arguments with multiple subobjects that +require special handling, and because C++17 does not support function templates +with multiple function parameter packs, an idiom is used to pass the kernel +arguments and special subobjects separately. +This idiom is best explained by way of an example. + +Consider the earlier example of a call to ``single_task()`` that passes a +lambda that captures variables of type ``int`` and ``std::stream``. +``std::stream`` is an example of a SYCL type that requires special handling in +kernel arguments. +The call to ``kernel_entry_point<KernelName>(kernelFunc)`` in the implementation +of ``single_task()`` results in an implicit call to ``sycl_kernel_launch`` that +looks similar to the following (the access to the captured copy of ``sout`` via +``kernelFunc.sout`` is not valid C++ syntax, but the compiler can generate such +accesses). + +.. code-block:: C++ + + sycl_kernel_launch<KernelName>("kernel-entry-point", kernelFunc)(kernelFunc.sout) + +The SYCL kernel name type, ``KernelName``, is passed as an explicit template +type argument for convenient use by the SYCL run-time library if desired. +The first function argument is the name of the offload kernel entry point +function generated for the SYCL kernel denoted by ``KernelName``. +This argument names the function that the SYCL run-time library, in conjunction +with an offload backend, needs to resolve and execute on the selected device +in order to execute the SYCL kernel. +The remaining arguments, for which there is just one in this case, +``kernelFunc``, are the kernel arguments. +Kernel arguments that contain a special subobject are passed as lvalues; those +that don't are passed as xvalues (and may therefore be moved from though such +types are unlikely to have optimized move operations since they are required to +be bit-copyable types). + +The result of the ``sycl_kernel_launch`` call must be a function object (e.g., +the result of a lambda expression, potentially one that captures references to +the kernel arguments). +The special subobjects are then passed as lvalues in an invocation of the +resulting function object. +This is the idiom mentioned earlier; this effectively allows for a call with +two function parameter packs. +A SYCL run-time library is expected to define the ``sycl_kernel_launch`` +template in a form similar to this: + +.. code-block:: C++ + + namespace sycl { + class handler { + template <typename KernelName, typename... KernelArgs> + auto sycl_kernel_launch(const char *entryPointName, KernelArgs &... args) { + return [&] (auto &... subobjects) { + // Process the kernel arguments and special subobjects, schedule + // execution of 'entryPointName' on the selected device, and + // return a type-list object sepcifying additional parameters to + // add to the offload kernel entry point function (see below). + return detail::type_list<...>{}; + }; + } + ... + }; + } + +The above call to a ``sycl_kernel_launch`` template suffices for Clang to inform +the SYCL run-time library of the name of the generated offload kernel entry +point function to be executed, the kernel arguments to be passed to it, and the +subobjects of those kernel arguments that require special handling. +This is all the information the SYCL run-time library needs from Clang to +fulfill its requirements. + +The special handling required for kernel arguments is an implementation detail +of the SYCL run-time library, but may require additional arguments to be passed +to the SYCL kernel. +This means that additional parameters may need to be added to the offload kernel +entry point function that Clang generates. +The SYCL run-time library informs Clang of the additional parameters by +returning a type-list object from the function object returned by the call to +``sycl_kernel_launch``. +A type-list object is a (possibly empty) object with a class template +specialization type; the template arguments for the specialization indicate +the additional parameters, if any, required for the kernel entry point. +The following class is suitable for type-list objects. + +.. code-block:: C++ + + namespace sycl { + namespace detail { + template <typename...> + class type_list {}; + } } - #endif -Depending on the compiler mode, ``multi_ptr`` will either decorate its internal -data with the address space attribute or not. +Continuing with the earlier example, assume that the ``sycl::stream`` type holds +a reference to a buffer and that the SYCL run-time implementation uses an +offload backend that requires the buffer to be passed as individual kernel +arguments of type ``buffer_t*`` and ``int``. +The implicit call to ``sycl_kernel_launch`` informed the SYCL run-time library +of the captured ``sycl::stream`` variable, ``sout``, by passing a reference to +the capture to the function object returned by ``sycl_kernel_launch``. +The SYCL run-time library is now obligated to report the additional required +parameters by returning an appropriate type-list from that function object. +In general, computing the type-list requires metaprogramming to inspect all of +the special subobject types. +In this example, that metaprogramming would ultimately result in an object of +type ``detail::type_list<buffer_t*, int>`` being returned. +The types of the template arguments of the type of the returned object direct +Clang to add additional parameters of type ``buffer_t*`` and ``int`` to the +kernel entry point it generates. -To utilize clang's existing functionality, we reuse the following OpenCL address -space attributes for pointers: +With the signature of the entry point function now known, there is just one +remaining bit of information that Clang needs to be informed of; what to do +with the additional parameters in the body of the generated kernel entry point. +The required handling of these is delegated to a +``sycl_handle_special_kernel_parameters`` template that is implicitly called +in similar fashion to the ``sycl_kernel_launch`` template. +The direct arguments are the references to the special subobjects of the +kernel parameters and ``sycl_handle_special_kernel_parameters`` is required +to return a function object with a call operator that can receive the +additional parameters. +For example: -.. list-table:: - :header-rows: 1 +.. code-block:: C++ + + namespace sycl { + class handler { + template <typename KernelName, typename... Subobjects> + static auto sycl_handle_special_kernel_parameters(Subobjects &... sos) { + return [&] (auto &... extraParams) -> void { + // Process the special subobjects in order consuming elements of + // 'extraParmas' as needed to update them. + }; + } + ... + }; + } + +With Clang now informed of the extra parameter handling necessitated by kernel +arguments with special subobjects, it has all the information needed to +generate the offload kernel entry point function. +The entry point function generated for the earlier example would look something +like the following (again, the use of ``kernelFunc.sout`` to access the captured +variable stored in ``kernelFunc`` is not valid C++ since captured variables +don't have names, but the intent should be clear; ``kernel-entry-point`` and +``lambda-from-f`` are exposition only names). + +.. code-block:: C++ + + void kernel-entry-point(lambda-from-f kernelFunc, buffer_t* X, int Y) { + sycl_handle_special_kernel_parameters(kernelFunc.sout)(X, Y); + kernelFunc(); + } + + +.. _attr-sycl_kernel_entry_point: + https://clang.llvm.org/docs/AttributeReference.html#sycl-kernel-entry-point +.. _SYCL-2020-3.13.1: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec::device.copyable +.. _SYCL-2020-4.9.4.2: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:invokingkernels +.. _SYCL-2020-4.12.4: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.parameter.passing +.. _SYCL-2020-5.2: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:naming.kernels + + +.. _sect-sycl_special_kernel_parameter: + +The ``[[clang::sycl_special_kernel_parameter]]`` Attribute +---------------------------------------------------------- +This attribute is intended for use in the implementation of SYCL run-time +libraries and should not be used directly by programmers. + +`SYCL 2020 section 4.12.4, "Rules for parameter passing to kernels" <SYCL-2020-4.12.4_>`_ +specifies that objects of some SYCL types may be passed as kernel arguments +(including as data members or captures of SYCL kernel objects) even if their +class definitions do not satisfy the device copyable requirements specified in +`SYCL 2020 section 3.13.1, "Device copyable" <SYCL-2020-3.13.1_>`_. +The SYCL run-time library provides the definition of these types and is thus +responsible for managing the details of how they are transferred to a device +and how they are received as parameters of the offload kernel entry point +function. + +Since C++17 lacks reflection features that would allow the SYCL run-time +library to identify use of these types for data members, captures, or base +classes of kernel argument types, the SYCL run-time library is dependent on +Clang to inform it of the presence of subobjects of these types within kernel +arguments. +The SYCL run-time library indicates which SYCL types require special handling +by declaring them with the +`sycl_special_kernel_parameter <attr-sycl_special_kernel_parameter_>`_ +attribute. + +When translating calls to functions declared with the +`sycl_kernel_entry_point <attr-sycl_kernel_entry_point_>`_ attribute, Clang +inspects each kernel argument type to look for data members, catures, and base +classes, that have a type declared with the +`sycl_special_kernel_parameter <attr-sycl_special_kernel_parameter_>`_ +attribute. +For each such occurrence, Clang passes a reference to the associated subobject +when constructing the implicit call to the ``sycl_kernel_launch`` template. +See the +:ref:`sycl_kernel_entry_point <sect-sycl_kernel_entry_point>` section for more +details. + +For example, kernel arguments that contain a subobject of +``sycl::local_accessor`` type usually cannot be bit-copied to device memory for +use by a SYCL kernel; additional support from the offload backend is generally +required to manage their associated memory. +A SYCL run-time library implementation might therefore declare this type +similar to the following. + +.. code-block:: C++ + + namespace sycl { + template <typename DataT, int Dimensions = 1> + class [[clang::sycl_special_kernel_parameter]] local_accessor { + ... + }; + } + +.. _attr-sycl_special_kernel_parameter: + https://clang.llvm.org/docs/AttributeReference.html#sycl-special-kernel-parameter + + +The ``[[clang::sycl_external]]`` Attribute +------------------------------------------ +This attribute is intended for use in the implementation of SYCL run-time +libraries and should not be used directly by programmers. - * - Address space attribute - - SYCL address_space enumeration - * - ``__attribute__((opencl_global))`` - - global_space, constant_space - * - ``__attribute__((opencl_global_device))`` - - global_space - * - ``__attribute__((opencl_global_host))`` - - global_space - * - ``__attribute__((opencl_local))`` - - local_space - * - ``__attribute__((opencl_private))`` - - private_space +The `sycl_external <attr-sycl_external_>`_ attribute implements the semantics +required for the ``SYCL_EXTERNAL`` macro specified in +`SYCL 2020 section 5.10.1, "SYCL functions and member functions linkage" <SYCL-2020-5.10.1_>`_. +A function defined with this attribute is emitted for each device target +regardless of whether it is ODR-used. +The ``SYCL_EXTERNAL`` macro should be defined by the ``<sycl/sycl.hpp>`` header +file or one that it includes. .. code-block:: C++ - //TODO: add support for __attribute__((opencl_global_host)) and __attribute__((opencl_global_device)). + #define SYCL_EXTERNAL [[clang::sycl_external]] +.. _attr-sycl_external: + https://clang.llvm.org/docs/AttributeReference.html#sycl-external +.. _SYCL-2020-5.10.1: + https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:syclexternal _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
