================
@@ -642,65 +654,87 @@ derived. Consider the following call to a SYCL kernel
invocation function.
});
}
-The SYCL kernel object is the result of the lambda expression. It has two
-data members corresponding to the captures of ``sout`` and ``s``. Since one
-of these data members corresponds to a special SYCL type that must be passed
-individually as an offload kernel parameter, it is necessary to decompose the
-SYCL kernel object into its constituent parts; the offload kernel will have
-two kernel parameters. Given a SYCL implementation that uses a
-``sycl_kernel_entry_point`` attributed function like the one shown above, an
-offload kernel entry point function will be generated that looks approximately
+The SYCL kernel object is the result of the lambda expression. The call to
+``kernel_entry_point()`` via the call to ``single_task()`` triggers the
+generation of an offload kernel entry point function that looks approximately
as follows.
.. code-block:: c++
- void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
- kernel-type kernel = { sout, s );
- kernel();
+ void sycl-kernel-caller-for-KN(kernel-type kernelFunc) {
+ kernelFunc();
}
There are a few items worthy of note:
-#. The name of the generated function incorporates the SYCL kernel name,
- ``KN``, that was passed as the ``KernelNameType`` template parameter to
- ``kernel_entry_point()`` and provided as the argument to the
- ``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
- between SYCL kernel names and offload kernel entry points.
+#. ``sycl-kernel-caller-for-KN`` is an exposition only name; the actual name
+ generated for an entry point is an implementation detail and subject to
+ change. However, the name will incorporate the SYCL kernel name, ``KN``,
+ that was passed as the ``KernelName`` template parameter to
+ ``single_task()`` and eventually provided as the argument to the
+ ``sycl_kernel_entry_point`` attribute in order to ensure that a unique
+ name is generated for each entry point. There is a one-to-one
correspondence
+ between SYCL kernel names and offload kernel entry points.
#. The SYCL kernel is a lambda closure type and therefore has no name;
``kernel-type`` is substituted above and corresponds to the ``KernelType``
- template parameter deduced in the call to ``kernel_entry_point()``.
- Lambda types cannot be declared and initialized using the aggregate
- initialization syntax used above, but the intended behavior should be clear.
+ template parameter deduced in the call to ``single_task()``.
+
+#. The parameter and the call to ``kernelFunc()`` in the function body
+ correspond to the definition of ``kernel_entry_point()`` as called by
+ ``single_task()``.
-#. ``S`` is a device copyable type that does not directly or indirectly contain
- a data member of a SYCL special type. It therefore does not need to be
- decomposed into its constituent members to be passed as a kernel argument.
+#. The parameter is type checked for conformance with the SYCL device
+ copyability and kernel parameter requirements.
-#. The depiction of the ``sycl::stream`` parameter as a single self contained
- kernel parameter is an oversimplification. SYCL special types may require
- additional decomposition such that the generated function might have three
- or more parameters depending on how the SYCL library implementation defines
- these types.
+Within ``single_task()``, the call to ``kernel_entry_point()`` is effectively
+replaced with a synthesized call to a ''sycl_kernel_launch`` template that
+looks approximately as follows.
-#. The call to ``kernel_entry_point()`` has no effect other than to trigger
- emission of the entry point function. The statments that make up the body
- of the function are not executed when the function is called; they are
- only used in the generation of the entry point function.
+.. code-block:: c++
+
+ sycl_kernel_launch<KN>("sycl-kernel-caller-for-KN", kernelFunc);
+
+There are a few items worthy of note:
+
+#. Lookup for the ``sycl_kernel_launch`` template is performed as if from the
+ body of the (possibly instantiated) definition of ``kernel_entry_point()``.
+ If name lookup or overload resolution fails, the program is ill-formed.
+ If the selected overload is a non-static member function, then ``this`` is
+ passed as the implicit object parameter.
+
+#. Function arguments passed to ``sycl_kernel_launch()`` are passed
+ as if by ``std::move(x)``.
+
+#. The ``sycl_kernel_launch`` template is expected to be provided by the SYCL
+ library implementation. It is responsible for copying the kernel arguments
+ to device memory and for scheduling execution of the generated offload
----------------
bader wrote:
Assuming that SYCL kernel type decomposition + compiler optimizations eliminate
the need for passing some kernel arguments how this information is supposed to
be communicated to the SYCL runtime library?
https://github.com/llvm/llvm-project/pull/152403
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
