[PATCH] D99190: WIP: [SYCL] Add design document for SYCL mode

Wed Mar 24 23:01:20 PDT 2021

keryell requested changes to this revision.
keryell added a comment.
This revision now requires changes to proceed.

Great to have some design documentation!
Just a feedback on the first 20%...

================
Comment at: clang/docs/SYCLSupport.md:23
+files (which are really LLVM IR files) are linked with the `llvm-link` tool.
+The resulting LLVM IR module is then translated into a SPIR-V module using the
+`llvm-spirv` tool and wrapped in a host object file using the
----------------
SPIR-V is just an example

================
Comment at: clang/docs/SYCLSupport.md:35
+- perform compilation separately from linkage
+- compile the device SPIR-V module ahead-of-time for one or more targets
+- perform device code splitting so that device code is distributed across
----------------

================
Comment at: clang/docs/SYCLSupport.md:48
+applies additional restrictions on the device code (e.g. no exceptions or
+virtual calls), generates LLVM IR for the device code only and "integration
+header" which provides information like kernel name, parameters order and data
----------------

================
Comment at: clang/docs/SYCLSupport.md:55
+  - Any LLVM IR transformation can be applied with only one limitation:
+    back-end compiler should be able to handle transformed LLVM IR. NOTE: loop
+    transformation passes performance impact depends on accurate target
----------------

================
Comment at: clang/docs/SYCLSupport.md:56
+    back-end compiler should be able to handle transformed LLVM IR. NOTE: loop
+    transformation passes performance impact depends on accurate target
+    information, so it make sense to disable such transformation for "virtual"
----------------

================
Comment at: clang/docs/SYCLSupport.md:57
+    transformation passes performance impact depends on accurate target
+    information, so it make sense to disable such transformation for "virtual"
+    targets like SPIR.
----------------

================
Comment at: clang/docs/SYCLSupport.md:60
+  - Optionally: Address space inference pass
+  - Optionally: LLVM IR → SPIR-V translator.
+- **Back-end** - produces native "device" code. It is shown as
----------------
Should we add somewhere other back-ends like PTX?

================
Comment at: clang/docs/SYCLSupport.md:69
+compiler to produce SYCL heterogeneous applications. Second, even if the
+same clang compiler is used for the host compilation, information provided in the
+integration header is used (included) by the SYCL runtime implementation, so the
----------------

================
Comment at: clang/docs/SYCLSupport.md:80
+- SYCL kernel function object (functor or lambda) lowering. This component
+creates an SPIR kernel function interface for SYCL kernels.
+- Device code diagnostics. This component enforces language restrictions on
----------------

================
Comment at: clang/docs/SYCLSupport.md:95
+...
+using namespace cl::sycl;
+queue Q;
----------------

================
Comment at: clang/docs/SYCLSupport.md:108
+In this example, the compiler needs to compile the lambda expression passed
+to the `cl::sycl::handler::parallel_for` method, as well as the function `foo`
+called from the lambda expression for the device.
----------------

================
Comment at: clang/docs/SYCLSupport.md:114
+portion of the source code (the contents of the lambda expression passed to the
+`cl::sycl::handler::parallel_for` and any function called from this lambda
+expression).
----------------

================
Comment at: clang/docs/SYCLSupport.md:118
+The current approach is to use the SYCL kernel attribute in the runtime to
+mark code passed to `cl::sycl::handler::parallel_for` as "kernel functions".
+The runtime library can't mark foo as "device" code - this is a compiler
----------------

================
Comment at: clang/docs/SYCLSupport.md:125
+
+All SYCL memory objects shared between host and device must be accessed through
+special `accessor` classes. The "device" side implementation of these classes
----------------
We need to introduce USM too somehow...

================
Comment at: clang/docs/SYCLSupport.md:127
+special `accessor` classes. The "device" side implementation of these classes
+contains pointers to the device memory. As there is no way in OpenCL to pass
+structures with pointers inside as kernel arguments all memory objects shared
----------------
To be generalized to non OpenCL case.

================
Comment at: clang/docs/SYCLSupport.md:139
+
+To facilitate the mapping of SYCL kernel data members to SPIR kernel arguments
+and overcome OpenCL limitations we added the generation of an SPIR kernel
----------------
To generalize to non SPIR & non OpenCL

================
Comment at: clang/docs/SYCLSupport.md:140
+To facilitate the mapping of SYCL kernel data members to SPIR kernel arguments
+and overcome OpenCL limitations we added the generation of an SPIR kernel
+function inside the compiler. A SPIR kernel function contains the body of the
----------------
It is not about overcoming some OpenCL limitations but just to use OpenCL as a back-end.

================
Comment at: clang/docs/SYCLSupport.md:155
+
+// Generated kernel function (expressed using OpenCL extensions)
+__kernel KernelName(global int* a) {
----------------
"OpenCL extensions" is probably misleading here and the sentence should be clarified.

================
Comment at: clang/docs/SYCLSupport.md:172
+document
+[SYCL Kernel Parameter Handling and Array Support](KernelParameterPassing.md).
+
----------------
Where is it?
I guess it is to be landed some day from https://github.com/intel/llvm/blob/sycl/sycl/doc/KernelParameterPassing.md ?

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D99190/new/

https://reviews.llvm.org/D99190