[SYCL][Doc] Add spec for "spirv" to kernel compiler #11954
Conversation
Add an extension that defines support for online compiled kernels that are written in binary SPIR-V.
Add a variant of `create_kernel_bundle_from_source` that reads the kernel source as a vector of bytes.
| |A C++ integral type with the same width. | ||
|
|
||
| |*OpTypeFloat* of width 32 | ||
| |`float` |
There was a problem hiding this comment.
Is there any reason why dtypes are mentioned here? As long as it is supported by SPIR-V, it should be fine.
There are other dtypes that are important for us and they are not mentioned i.e. bfloat16 or float16/half
There was a problem hiding this comment.
In DPC++, the implementations of both half and bfloat16 are actually a struct around a 16-bit integer. Therefore, I think you would pass parameters of these types as OpTypeInt of width 16. Inside the kernel, you would treat the value as either "half" or "bfloat16".
There was a problem hiding this comment.
I loosened the wording here in af5e8c0.
With the new wording, it should be possible to pass sycl::half for a SPIR-V 16-bit OpTypeFloat parameter. It should also be possible to pass sycl::ext::oneapi::bfloat16 for a SPIR-V "bf16 type" (which is actually represented in SPIR-V as a 16-bit OpTypeInt).
The kernel compiler no longer supports a mechanism to get a kernel by its type-name, so remove this reference about the type-name.
Be a little less proscriptive when describing how to pass kernel arguments.
The SPIR-V widths are in bits not bytes.
sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc
Outdated
Show resolved
Hide resolved
|
Maybe this is opening too big a can of worms - but it's a little odd that we're calling binary SPIRV "source". In OpenCL there is createProgramFromSource AND createProgramFromIL. |
It's not clear what it would mean to create a kernel bundle from raw bytes if the kernel was written in a textual language (e.g. OpenCL). It seems like we would need to define the character encoding, and that seems like a bad idea. Instead, classify each source language as either text format or binary format. We now require the application to use the overload of `create_kernel_bundle_from_source` that matches the format.
I'm, standing firm on this one. I think the important distinction is whether the language is text format or binary format, and we can deal with this nicely by having two overloads of I also think the distinction of "source" and "IL" could get fuzzy. For example, what if we want to allow kernels to be defined using the textual form of SPIR-V in the future? Would we call this "source" (and not "IL") just because we need to run the program through the SPIR-V assembler? I think it's better to expose just a single function name |
|
@intel/llvm-gatekeepers I think this is ready to merge. I talked with @jbrodman offline, and he does not feel so strongly about his comment that he opposes this PR. |
Implements the extension described in intel#11954. This PR includes the following changes: - Adds a `create_kernel_bundle_from_source` overload for `std::vector<std::byte>` kernel sources. - Adds new `source_language::spirv`. - Defines `SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV`. - Adds support for SPIR-V kernels created from `std::vector<std::byte>` sources. - Moves `sycl_ext_oneapi_kernel_compiler_spirv.asciidoc` from `proposed` to `experimental`. Signed-off-by: Michael Aziz <[email protected]>
Implements the extension described in #11954. This PR includes the following changes: - Adds a `create_kernel_bundle_from_source` overload for `std::vector<std::byte>` kernel sources. - Adds new `source_language::spirv`. - Defines `SYCL_EXT_ONEAPI_KERNEL_COMPILER_SPIRV`. - Adds support for SPIR-V kernels created from `std::vector<std::byte>` sources. - Moves `sycl_ext_oneapi_kernel_compiler_spirv.asciidoc` from `proposed` to `experimental`. --------- Signed-off-by: Michael Aziz <[email protected]>
Add an extension that defines support for online compiled kernels that
are written in binary SPIR-V.