[clang] [SYCL] SYCL host kernel launch support for the sycl_kernel_entry_point attribute. (PR #152403)

Tom Honermann via cfe-commits cfe-commits at lists.llvm.org
Thu Nov 20 22:14:52 PST 2025


================
@@ -683,24 +696,55 @@ There are a few items worthy of note:
    or more parameters depending on how the SYCL library implementation defines
    these types.
 
-#. 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.
+The call to ``kernel_entry_point()`` by ``single_task()`` is effectively
+replaced with synthesized code that looks approximately as follows.
+
+.. code-block:: c++
+
+  sycl::stream sout = Kernel.sout;
----------------
tahonermann wrote:

Thank you for point out that the use of `Kernel` here wasn't particularly clear. When I went and read back through the doc, I found quite a bit I wasn't very happy with, so I did quite a bit of rewriting today. I hope it is better than it was.

I kept use of `sycl::stream` in the example for now because it is conceptually easy to understand; it evokes thoughts of hello world generally. However, I took out all of the mention of kernel argument decomposition and reconstruction since such support hasn't been implemented yet. The stream class is barely mentioned now, so hopefully that helps. Use of USM would make for a simple example, but it hides the host/device distinction that I think is important to understand. Use of buffers and accessors would be ok, but makes the example more complicated. Maybe we can revisit this as part of adding support for SYCL special types; that might require a more realistic example anyway.

https://github.com/llvm/llvm-project/pull/152403


More information about the cfe-commits mailing list