@@ -455,6 +455,180 @@ The SYCL kernel in the previous code sample meets these expectations.
455455 }];
456456}
457457
458+ def SYCLKernelEntryPointDocs : Documentation {
459+ let Category = DocCatFunction;
460+ let Content = [{
461+ The ``sycl_kernel_entry_point`` attribute facilitates the generation of an
462+ offload kernel entry point, sometimes called a SYCL kernel caller function,
463+ suitable for invoking a SYCL kernel on an offload device. The attribute is
464+ intended for use in the implementation of SYCL kernel invocation functions
465+ like the ``single_task`` and ``parallel_for`` member functions of the
466+ ``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
467+ class", of the SYCL 2020 specification.
468+
469+ The attribute requires a single type argument that specifies a class type that
470+ meets the requirements for a SYCL kernel name as described in section 5.2,
471+ "Naming of kernels", of the SYCL 2020 specification. A unique kernel name type
472+ is required for each function declared with the attribute. The attribute may
473+ not first appear on a declaration that follows a definition of the function.
474+
475+ The attribute only appertains to functions and only those that meet the
476+ following requirements.
477+
478+ * Has a ``void`` return type.
479+ * Is not a non-static member function, constructor, or destructor.
480+ * Is not a C variadic function.
481+ * Is not a coroutine.
482+ * Is not defined as deleted or as defaulted.
483+ * Is not declared with the ``constexpr`` or ``consteval`` specifiers.
484+ * Is not declared with the ``[[noreturn]]`` attribute.
485+
486+ Use in the implementation of a SYCL kernel invocation function might look as
487+ follows.
488+
489+ .. code-block:: c++
490+
491+ namespace sycl {
492+ class handler {
493+ template<typename KernelNameType, typename KernelType>
494+ [[ clang::sycl_kernel_entry_point(KernelNameType) ]]
495+ static void kernel_entry_point(KernelType kernel) {
496+ kernel();
497+ }
498+
499+ public:
500+ template<typename KernelNameType, typename KernelType>
501+ void single_task(KernelType kernel) {
502+ // Call kernel_entry_point() to trigger generation of an offload
503+ // kernel entry point.
504+ kernel_entry_point<KernelNameType>(kernel);
505+ // Call functions appropriate for the desired offload backend
506+ // (OpenCL, CUDA, HIP, Level Zero, etc...).
507+ }
508+ };
509+ } // namespace sycl
510+
511+ A SYCL kernel is a callable object of class type that is constructed on a host,
512+ often via a lambda expression, and then passed to a SYCL kernel invocation
513+ function to be executed on an offload device. A SYCL kernel invocation function
514+ is responsible for copying the provided SYCL kernel object to an offload
515+ device and initiating a call to it. The SYCL kernel object and its data members
516+ constitute the parameters of an offload kernel.
517+
518+ A SYCL kernel type is required to satisfy the device copyability requirements
519+ specified in section 3.13.1, "Device copyable", of the SYCL 2020 specification.
520+ Additionally, any data members of the kernel object type are required to satisfy
521+ section 4.12.4, "Rules for parameter passing to kernels". For most types, these
522+ rules require that the type is trivially copyable. However, the SYCL
523+ specification mandates that certain special SYCL types, such as
524+ ``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are not
525+ trivially copyable. These types require special handling because they cannot
526+ be copied to device memory as if by ``memcpy()``. Additionally, some offload
527+ backends, OpenCL for example, require objects of some of these types to be
528+ passed as individual arguments to the offload kernel.
529+
530+ An offload kernel consists of an entry point function that declares the
531+ parameters of the offload kernel and the set of all functions and variables that
532+ are directly or indirectly used by the entry point function.
533+
534+ A SYCL kernel invocation function invokes a SYCL kernel on a device by
535+ performing the following tasks (likely with the help of an offload backend
536+ like OpenCL):
537+
538+ #. Identifying the offload kernel entry point to be used for the SYCL kernel.
539+
540+ #. Deconstructing the SYCL kernel object, if necessary, to produce the set of
541+ offload kernel arguments required by the offload kernel entry point.
542+
543+ #. Copying the offload kernel arguments to device memory.
544+
545+ #. Initiating execution of the offload kernel entry point.
546+
547+ The offload kernel entry point for a SYCL kernel performs the following tasks:
548+
549+ #. Reconstituting the SYCL kernel object, if necessary, using the offload
550+ kernel parameters.
551+
552+ #. Calling the ``operator()`` member function of the (reconstituted) SYCL kernel
553+ object.
554+
555+ The ``sycl_kernel_entry_point`` attribute automates generation of an offload
556+ kernel entry point that performs those latter tasks. The parameters and body of
557+ a function declared with the ``sycl_kernel_entry_point`` attribute specify a
558+ pattern from which the parameters and body of the entry point function are
559+ derived. Consider the following call to a SYCL kernel invocation function.
560+
561+ .. code-block:: c++
562+
563+ struct S { int i; };
564+ void f(sycl::handler &handler, sycl::stream &sout, S s) {
565+ handler.single_task<struct KN>([=] {
566+ sout << "The value of s.i is " << s.i << "\n";
567+ });
568+ }
569+
570+ The SYCL kernel object is the result of the lambda expression. It has two
571+ data members corresponding to the captures of ``sout`` and ``s``. Since one
572+ of these data members corresponds to a special SYCL type that must be passed
573+ individually as an offload kernel parameter, it is necessary to decompose the
574+ SYCL kernel object into its constituent parts; the offload kernel will have
575+ two kernel parameters. Given a SYCL implementation that uses a
576+ ``sycl_kernel_entry_point`` attributed function like the one shown above, an
577+ offload kernel entry point function will be generated that looks approximately
578+ as follows.
579+
580+ .. code-block:: c++
581+
582+ void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
583+ kernel-type kernel = { sout, s );
584+ kernel();
585+ }
586+
587+ There are a few items worthy of note:
588+
589+ #. The name of the generated function incorporates the SYCL kernel name,
590+ ``KN``, that was passed as the ``KernelNameType`` template parameter to
591+ ``kernel_entry_point()`` and provided as the argument to the
592+ ``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
593+ between SYCL kernel names and offload kernel entry points.
594+
595+ #. The SYCL kernel is a lambda closure type and therefore has no name;
596+ ``kernel-type`` is substituted above and corresponds to the ``KernelType``
597+ template parameter deduced in the call to ``kernel_entry_point()``.
598+ Lambda types cannot be declared and initialized using the aggregate
599+ initialization syntax used above, but the intended behavior should be clear.
600+
601+ #. ``S`` is a device copyable type that does not directly or indirectly contain
602+ a data member of a SYCL special type. It therefore does not need to be
603+ decomposed into its constituent members to be passed as a kernel argument.
604+
605+ #. The depiction of the ``sycl::stream`` parameter as a single self contained
606+ kernel parameter is an oversimplification. SYCL special types may require
607+ additional decomposition such that the generated function might have three
608+ or more parameters depending on how the SYCL library implementation defines
609+ these types.
610+
611+ #. The call to ``kernel_entry_point()`` has no effect other than to trigger
612+ emission of the entry point function. The statments that make up the body
613+ of the function are not executed when the function is called; they are
614+ only used in the generation of the entry point function.
615+
616+ It is not necessary for a function declared with the ``sycl_kernel_entry_point``
617+ attribute to be called for the offload kernel entry point to be emitted. For
618+ inline functions and function templates, any ODR-use will suffice. For other
619+ functions, an ODR-use is not required; the offload kernel entry point will be
620+ emitted if the function is defined.
621+
622+ Functions declared with the ``sycl_kernel_entry_point`` attribute are not
623+ limited to the simple example shown above. They may have additional template
624+ parameters, declare additional function parameters, and have complex control
625+ flow in the function body. Function parameter decomposition and reconstitution
626+ is performed for all function parameters. The function must abide by the
627+ language feature restrictions described in section 5.4, "Language restrictions
628+ for device functions" in the SYCL 2020 specification.
629+ }];
630+ }
631+
458632def SYCLSpecialClassDocs : Documentation {
459633 let Category = DocCatStmt;
460634 let Content = [{
0 commit comments