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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
