diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index 21c7719a50b91..d796e49b628a7 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -10,6 +10,7 @@ :encoding: utf-8 :lang: en :dpcpp: pass:[DPC++] +:cpp: pass:[C++] :endnote: —{nbsp}end{nbsp}note // Set the default source code type in this document to C++, @@ -45,15 +46,17 @@ This extension also depends on the following other SYCL extensions: * link:../experimental/sycl_ext_oneapi_properties.asciidoc[ sycl_ext_oneapi_properties] +* link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[ + sycl_ext_oneapi_free_function_kernels] == Status -This is an experimental extension specification, intended to provide early -access to features and gather community feedback. Interfaces defined in -this specification are implemented in DPC++, but they are not finalized -and may change incompatibly in future versions of DPC++ without prior notice. -*Shipping software products should not rely on APIs defined in +This is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in +this specification are implemented in {dpcpp}, but they are not finalized +and may change incompatibly in future versions of {dpcpp} without prior notice. +*Shipping software products should not rely on APIs defined in this specification.* @@ -61,11 +64,12 @@ this specification.* This extension adds APIs that allow the application to dynamically generate the source code for a kernel, which it can then compile and enqueue to a device. -The APIs allow a kernel to be written in any of several possible languages, -where support for each of these languages is defined by a separate extension. -As a result, this extension provides a framework of APIs for online compilation -of kernels, but it does not define the details for any specific kernel language. -These details are provided by other extensions. +This extension provides support for kernels written in SYCL according to the +"free function kernel" syntax defined in +link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[ +sycl_ext_oneapi_free_function_kernels]. +However, other extensions could provide support for writing kernels in other +languages. The new APIs added by this extension are an expansion of the existing `kernel_bundle` capabilities. @@ -105,7 +109,8 @@ supports. This extension adds the `ext_oneapi_source` enumerator to `sycl::bundle_state` to identify a kernel bundle that is represented as a source code string. -``` +[source,c++] +---- namespace sycl { enum class bundle_state : /*unspecified*/ { @@ -114,26 +119,30 @@ enum class bundle_state : /*unspecified*/ { }; } // namespace sycl -``` +---- === New enumerator of kernel source languages This extension adds the `source_language` enumeration, which identifies possible languages for a kernel bundle that is in `ext_oneapi_source` state: -``` +[source,c++] +---- namespace sycl::ext::oneapi::experimental { enum class source_language : /*unspecified*/ { - // see below + sycl }; } // namespace sycl::ext::oneapi::experimental -``` +---- -However, there are no enumerators defined by this extension. -Instead, this enumeration is an extension point for other extensions, which can -specify the exact semantics of each possible kernel language. +The only enumerator defined by this extension is `sycl`, which indicates that +the kernel is written in SYCL using the "free function kernel" syntax defined +in link:../proposed/sycl_ext_oneapi_free_function_kernels.asciidoc[ +sycl_ext_oneapi_free_function_kernels]. +Other extensions may provide other enumerators that correspond to other +languages. === New member functions for the device class @@ -144,7 +153,7 @@ a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- class device { @@ -168,30 +177,32 @@ a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- namespace sycl::ext::oneapi::experimental { +template (1) kernel_bundle create_kernel_bundle_from_source( const context& ctxt, source_language lang, - const std::string& source) + const std::string& source, + PropertyListT props = {}) +template (2) kernel_bundle create_kernel_bundle_from_source( const context& ctxt, source_language lang, - const std::vector& bytes) + const std::vector& bytes, + PropertyListT props = {}) } // namespace sycl::ext::oneapi::experimental ---- !==== -_Preconditions:_ There are two overloads of this function: one that reads the -source code of the kernel from a `std::string`, and one that reads the source -code of the kernel from a `std::vector` of `std::byte`. -Each source language `lang` specifies whether the language is text format or -binary format, and the application must use the overload that corresponds to -that format. +_Constraints:_ Available only when `PropertyListT` is an instance of +`sycl::ext::oneapi::experimental::properties` which contains no properties +other than those listed below in the section "New properties for the +`create_kernel_bundle_from_source` function". _Effects:_ Creates a new kernel bundle that represents a kernel written in the source language `lang`, where the source code is contained either by `source` @@ -202,6 +213,14 @@ may only be submitted to a queue that shares the same context. The bundle's set of associated devices is the set of devices contained in `ctxt`. +Each source language `lang` specifies whether the language is text format or +binary format, and the application must use the overload that corresponds to +that format. +Applications must use overload (1) when the source language is text format and +must use overload (2) when the source language is binary format. +The `sycl` language is text format, so application must use overload (1) when +creating a kernel bundle from this language. + _Returns:_ The newly created kernel bundle, which has `ext_oneapi_source` state. @@ -209,10 +228,16 @@ _Throws:_ * An `exception` with the `errc::invalid` error code if the source language `lang` is not supported by any device contained by the context `ctxt`. +* An `exception` with the `errc::invalid` error code if the source language + `lang` does not support one of the properties in `PropertyListT`. +* Overload (1) throws an `exception` with the `errc::invalid` error code if the + source language `lang` is binary format. +* Overload (2) throws an `exception` with the `errc::invalid` error code if the + source language `lang` is text format. [_Note:_ Calling this function does not attempt to compile the source code. -As a result, syntactic errors in the source code string are not diagnosed by -this function. +As a result, syntax errors in `source` or `bytes` are not diagnosed by this +function. This function succeeds even if some devices in `ctxt` do not support the source language `lang`. @@ -225,7 +250,7 @@ a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- namespace sycl::ext::oneapi::experimental { @@ -246,7 +271,8 @@ kernel_bundle build( _Constraints:_ Available only when `PropertyListT` is an instance of `sycl::ext::oneapi::experimental::properties` which contains no properties -other than those listed below in the section "New properties". +other than those listed below in the section "New properties for the `build` +function". _Effects (1):_ The source code from `sourceBundle` is translated into one or more device images of state `bundle_state::executable`, and a new kernel bundle is @@ -272,8 +298,10 @@ _Throws:_ `devs` does not support compilation of kernels in the source language of `sourceBundle`. -* An `exception` with the `errc::invalid` error code if `props` contains an - `options` property that specifies an invalid option. +* An `exception` with the `errc::invalid` error code if the source language + `lang` does not support one of the properties in `PropertyListT` or if + `props` contains a `build_options` property that contains an option that is + not supported by `lang`. * An `exception` with the `errc::build` error code if the compilation or linking operations fail. @@ -290,7 +318,81 @@ exception is caught and handled appropriately. _{endnote}_] |==== -=== New properties +=== New properties for the `create_kernel_bundle_from_source` function + +This extension adds the following properties, which can be used in conjunction +with the `create_kernel_bundle_from_source` function that is defined above: + +|==== +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +struct include_files { + include_files(); (1) + include_files(const std::string &name, const std::string &content); (2) + void add(const std::string &name, const std::string &content); (3) +}; +using include_files_key = include_files; + +template<> +struct is_property_key : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +---- +!==== + +This property provides the name and content of include files that can be +referenced from the source code in the `source` parameter to +`create_kernel_bundle_from_source`. +The property conceptually contains a collection of (_Name_, _Content_) pairs, +where both _Name_ and _Content_ are strings. +The _Name_ is the name of an include file and the _Content_ is the content of +that include file. + +When the source language is `source_language::sycl`, the source code can have +`#include` statements where the name and content of the include file is +defined by this property. +For example, if the source code has `#include "foo/bar.h"`, the compilation +process will look at the `include_files` property to see if there is an entry +whose _Name_ is `foo/bar.h`. +If such an entry is found, the compiler uses the associated _Content_ as the +content of the include file. + +When the source language is `source_language::sycl`, the following header files +are implicitly available. +Therefore, the source string may `#include` these even without defining their +content via the `include_files` property: + +* ``; +* The {cpp} standard library headers; +* The SYCL backend headers `"sycl/backend/.hpp"` for any backends + that the implementation supports; and +* Any SYCL extension headers in "sycl/ext" for extensions that the + implementation supports. + +The include files defined via the `include_files` property are searched first, +before these implicitly available headers. + +_Effects (1):_ Creates a new `include_files` property with no (_Name_, +_Content_) pairs. + +_Effects (2):_ Creates a new `include_files` property with a single (_Name_, +_Content_) pair. + +_Effects (3):_ Adds a (_Name_, _Content_) pair to the property. + +_Throws (3):_ + +* An `exception` with the `errc::invalid` error code if there is already an + entry with `name` in this property. +|==== + +=== New properties for the `build` function This extension adds the following properties, which can be used in conjunction with the `build` function that is defined above: @@ -300,14 +402,15 @@ a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- namespace sycl::ext::oneapi::experimental { struct build_options { - std::vector opts; - build_options(const std::string &opt); (1) - build_options(const std::vector &opts); (2) + build_options(); (1) + build_options(const std::string &opt); (2) + build_options(const std::vector &opts); (3) + void add(const std::string &opt); (4) }; using build_options_key = build_options; @@ -317,27 +420,30 @@ using build_options_key = build_options; This property provides build options that may affect the compilation or linking of the kernel, where each build option is a string. -There are no standard build options that are common across all source -languages. -Instead, each source language specification defines its own set of build -options. +All source languages support the `build_options` property, but each source +language defines the specific options that it supports. +The `source_language::sycl` language does not define any standard build +options, but an implementation may support implementation-defined options. -_Effects (1):_ Constructs a `build_options` property with a single build +_Effects (1):_ Constructs a `build_options` property with no build options. + +_Effects (2):_ Constructs a `build_options` property with a single build option. -_Effects (2):_ Constructs a `build_options` property from a vector of build +_Effects (3):_ Constructs a `build_options` property from a vector of build options. +_Effects (4):_ Adds `opt` to the end of the property's list of build options. + a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- namespace sycl::ext::oneapi::experimental { struct save_log { - std::string *log; save_log(std::string *to); (1) }; using save_log_key = save_log; @@ -356,15 +462,77 @@ provides a description of the error. Instead, the `save_log` property provides information about a build operation that succeeds. This might include warning messages or other diagnostics. -Each source language specification can define specific information that is -provided in the log. +All source languages support the `save_log` property, but each source language +defines the specific information that is provided in the log. +The `source_language::sycl` language does not define any specific information +that is provided in the log, so implementations are free to provide any +information they choose here. In general, the log information is intended for human consumption, and the format may not be stable across implementations of this extension. _Effects (1):_ Constructs a `save_log` property with a pointer to a `std::string`. -When the `build` function completes successfully, this string will contain the -log. +If the `to` pointer is not null, when the `build` function completes +successfully, the string pointed at by `to` will contain the log. + +_Remarks (1):_ When `to` is not null, the string object it points to must +remain valid for all calls to `build` taking this `save_log` property. + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +struct registered_kernel_names { + registered_kernel_names(); (1) + registered_kernel_names(const std::string &name); (2) + registered_kernel_names(const std::vector &names); (3) + void add(const std::string &name); (4) +}; +using registered_kernel_names_key = registered_kernel_names; + +template<> +struct is_property_key : std::true_type {}; + +} // namespace sycl::ext::oneapi::experimental +---- +!==== +This property is useful when the source language represents kernel names +differently in the source code and the generated code. +For example, {cpp} function names in the generated code are "mangled" in an +implementation-defined way. +The precise meaning of this property is defined by each source language, but in +general it allows the application to supply a list of kernel names as they +appear in the source code. +The application can then get the corresponding raw (i.e. mangled) names after +the code is compiled. +See the section below "Obtaining a kernel when the language is ``sycl``" for a +description of how this property is used with the `source_language::sycl` +language. + +_Effects (1):_ Creates a new `registered_kernel_names` property with no +registered kernel names. + +_Effects (2):_ Creates a new `registered_kernel_names` property with a single +registered kernel name. + +_Effects (3):_ Creates a new `registered_kernel_names` property from a vector +of kernel names. + +_Effects (4):_ Adds `name` to the property's list of registered kernel names. + +_Preconditions (2-4):_ Each source language defines its own requirements for +the registered kernel names. +For the language `source_language::sycl`, each name must be a {cpp} expression +for a pointer to a kernel function as defined below under "Obtaining a kernel +when the language is ``sycl``". + +[_Note:_ It is not an error to have duplicate names in a +`registered_kernel_names` property, but the duplicates have no effect. +_{endnote}_] |==== === New constraint for kernel bundle member functions @@ -403,7 +571,8 @@ of `kernel_id` objects if the kernel bundle was created from a bundle of state This extensions adds the following new `kernel_bundle` member functions: -``` +[source,c++] +---- namespace sycl { template @@ -412,17 +581,18 @@ class kernel_bundle { bool ext_oneapi_has_kernel(const std::string &name); kernel ext_oneapi_get_kernel(const std::string &name); + std::string ext_oneapi_get_raw_kernel_name(const std::string &name); }; } // namespace sycl -``` +---- |==== a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- bool ext_oneapi_has_kernel(const std::string &name) ---- @@ -441,7 +611,7 @@ a| [frame=all,grid=none] !==== a! -[source] +[source,c++] ---- kernel ext_oneapi_get_kernel(const std::string &name) ---- @@ -455,51 +625,316 @@ is `name`. _Throws:_ +* An `exception` with the `errc::invalid` error code if + `ext_oneapi_has_kernel(name)` returns `false`. + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +std::string ext_oneapi_get_raw_kernel_name(const std::string &name) +---- +!==== + +_Constraints:_ This function is not available when `State` is +`bundle_state::ext_oneapi_source`. + +_Returns:_ If the kernel bundle was created from a bundle of state +`bundle_state::ext_oneapi_source` and `name` was registered via +`registered_kernel_names`, returns the compiler-generated (e.g. mangled) name +for this kernel function. +If the kernel bundle was created from a bundle of state +`bundle_state::ext_oneapi_source` and `name` is the same as a +compiler-generated name for a kernel defined in that bundle, that same +`name` is returned. + +_Throws:_ + * An `exception` with the `errc::invalid` error code if `ext_oneapi_has_kernel(name)` returns `false`. |==== +=== Obtaining a kernel when the language is `sycl` + +When the kernel is defined in the language `source_language::sycl`, the host +code may query for the kernel or obtain the `kernel` object using either the +kernel's name as it is generated by the compiler (i.e. the {cpp} mangled name) +or by using the `registered_kernel_names` property. + +==== Using the compiler-generated name + +If the kernel is declared as `extern "C"`, the compiler generates the kernel +name exactly as it appears in the source code (i.e. there is no name mangling). +Therefore, it is easy to query for the kernel by using the compiler-generated +name. +For example, if the kernel is defined like this in the source code string: + +[source,c++] +---- +std::string source = R"""( + #include + namespace syclexp = sycl::ext::oneapi::experimental; + + extern "C" + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void foo(int *in, int *out) {/*...*/} +)"""; +---- + +Then the application's host code can query for the kernel like this: + +[source,c++] +---- +sycl::kernel_bundle kb = /*...*/; +sycl::kernel k = kb.ext_oneapi_get_kernel("foo"); +---- + +==== Using the `registered_kernel_names` property + +When the kernel is not declared as `extern "C"`, the compiler generates a +mangled name, so it is more convenient to use the `registered_kernel_names` +property. +Each string in the property must be the {cpp} expression for a pointer to a +kernel function. +These expression strings are conceptually compiled at the bottom of source +code. +To illustrate, consider source code that defines a kernel like this: + +[source,c++] +---- +std::string source = R"""( + #include + namespace syclexp = sycl::ext::oneapi::experimental; + + namespace mykernels { + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void bar(int *in, int *out) {/*...*/} + + } // namespace mykernels +)"""; +---- + +The host code can compile this and get the kernel's `kernel` object like so: + +[source,c++] +---- +sycl::kernel_bundle kb_src = /*...*/; + +sycl::kernel_bundle kb = syclexp::build(kb_src, + syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); -== Example +sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar"); +---- + +The {cpp} expression `"mykernels::bar"` computes the address of the kernel +function `bar`. +The host code then passes the same string (`"mykernels::bar"`) to +`ext_oneapi_get_kernel` in order to get the `kernel` object. +The string must have exactly the same content as the string that was used to +construct the property, without even any whitespace differences. + +The application can also obtain the compiler-generated (i.e. mangled) name for +the kernel by calling `ext_oneapi_get_raw_kernel_name` like this: + +[source,c++] +---- +sycl::kernel_bundle kb_src = /*...*/; + +sycl::kernel_bundle kb = syclexp::build(kb_src, + syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); + +std::string mangled_name = kb.ext_oneapi_get_raw_kernel_name("mykernels::bar"); +---- + +Again, the string passed to `ext_oneapi_get_raw_kernel_name` must have exactly +the same content as the string that was used to construct the +`registered_kernel_names` property. +The application may also pass this compiler-generated (i.e. mangled) name to +`ext_oneapi_get_kernel` in order to get the `kernel` object. + +==== Instantiating templated kernel functions + +The `registered_kernel_names` property can also be used to instantiate a +kernel that is defined as a function template. +For example, consider source code that defines a kernel function template like +this: + +[source,c++] +---- +std::string source = R"""( + #include + namespace syclexp = sycl::ext::oneapi::experimental; + + template + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void bartmpl(T *in, T *out) {/*...*/} +)"""; +---- + +The application can use the `registered_kernel_names` property to instantiate +the template for specific template arguments. +For example, this host code instantiates the template twice and gets a `kernel` +object for each instantiation: + +[source,c++] +---- +sycl::kernel_bundle kb_src = /*...*/; + +sycl::kernel_bundle kb = syclexp::build(kb_src, + syclexp::properties{syclexp::registered_kernel_names{{"bartmpl", "bartmpl"}}); + +sycl::kernel k_float = kb.ext_oneapi_get_kernel("bartmpl"); +sycl::kernel k_int = kb.ext_oneapi_get_kernel("bartmpl"); +---- + + +== Examples + +=== Simple example The following example demonstrates how a SYCL application can define a kernel as a string and then compile and launch it. -``` +[source,c++] +---- #include -namespace syclex = sycl::ext::oneapi::experimental; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; int main() { sycl::queue q; - // The source code for one or more kernels, defined in one of - // the supported source languages. + // The source code for a kernel, defined as a SYCL "free function kernel". std::string source = R"""( - /* language specific kernel source code */ + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + extern "C" + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } )"""; - // Create a kernel bundle in "source" state. The "some-language" is - // a stand-in for the enumerator telling which source language is used. + // Create a kernel bundle in "source" state. sycl::kernel_bundle kb_src = - syclex::create_kernel_bundle_from_source( + syclexp::create_kernel_bundle_from_source( q.get_context(), - syclex::source_language::/*some-language*/, + syclexp::source_language::sycl, source); + // Compile the kernel. There is no need to use the "registered_kernel_names" + // property because the kernel is declared extern "C". sycl::kernel_bundle kb_exe = - syclex::build(kb_src); + syclexp::build(kb_src); - // Get the kernel via its name. The "kernel-name" is a stand-in for the - // actual kernel name in the source string. - sycl::kernel k = kb_exe.ext_oneapi_get_kernel("kernel-name"); + // Get the kernel via its compiler-generated name. + sycl::kernel iota = kb_exe.ext_oneapi_get_kernel("iota"); + float *ptr = sycl::malloc_shared(NUM, q); q.submit([&](sycl::handler &cgh) { - // Any arguments for the kernel must be set manually. - cgh.set_args(/*...*/); + // Set the values of the kernel arguments. + cgh.set_args(3.14f, ptr); - // Launch the kernel according to its type. - // This assumes a simple "range" kernel. - cgh.parallel_for(sycl::range{1024}, k); - }); + // Launch the kernel according to its type, in this case an nd-range kernel. + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, iota); + }).wait(); } -``` +---- + +=== Disambiguating overloaded kernel functions + +This example demonstrates how to use the `registered_kernel_names` property to +disambiguate a kernel function that has several overloads. + +[source,c++] +---- +#include +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +int main() { + sycl::queue q; + + // The source code for two kernels defined as overloaded functions. + std::string source = R"""( + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::range_kernel<1>)) + void iota(int start, int *ptr) { + size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } + )"""; + + // Create a kernel bundle in "source" state. + sycl::kernel_bundle kb_src = + syclexp::create_kernel_bundle_from_source( + q.get_context(), + syclexp::source_language::sycl, + source); + + // Compile the kernel. Because there are two overloads of "iota", we need to + // use a C++ cast to disambiguate between them. Here, we are selecting the + // "int" overload. + std::string iota_name{"(void(*)(int, int*))iota"}; + sycl::kernel_bundle kb_exe = syclexp::build(kb_src, + syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + + // Get the kernel by passing the same string we used to construct the + // "registered_kernel_names" property. + sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); + + int *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + // Set the values of the kernel arguments. + cgh.set_args(3, ptr); + + // Launch the kernel according to its type, in this case an nd-range kernel. + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, iota); + }).wait(); +} +---- + + +== Issues + +* Do we want to add an API similar to `nvrtcGetTypeName`? + This does seem useful in some advanced cases, and it is not specific to CUDA. + The implementation is fairly straightforward. + You use `typeid` to get a `std::type_info`. + You can then call `type_info::name` to get an implementation-defined name for + the type. + For clang on Linux, this returns the type's mangled name. + You can then call `+abi::__cxa_demangle+` to get an unmangled name for the + type. + I'm not sure about the details on Windows hosts, though. + If `type_info::name` returns a mangled name on Windows too, then maybe we can + still use `+abi::__cxa_demangle+` to get an unmangled name, but this needs to + be checked. ++ +Another option might be to provide this functionality as a utility library. +There is no inherent reason why this functionality needs to be built into +{dpcpp}. +However, we don't yet have a utility library where this would go, and it may be +hard for customers to discover this functionality if it is defined outside of +this extension. diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index 51ea0fc5d72c4..7d79e75634c4d 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -254,11 +254,15 @@ If the device does not support kernel bundles written in |==== -=== Build options +=== Supported properties -The `build_options` property accepts any of the compiler or linker options -defined by the OpenCL specification, except for those that are specific to -creating an OpenCL library. +The `opencl` language supports only those properties from +link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[ +sycl_ext_oneapi_kernel_compiler] that are common to all source languages. + +The `build_options` property that is passed to the `build` function may contain +any of the compiler or linker options defined by the OpenCL specification, +except for those that are specific to creating an OpenCL library. The kernel compiler can be used to create an OpenCL program, but not an OpenCL library. diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc index ca89d4b744879..c985c5b149274 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_spirv.asciidoc @@ -123,11 +123,15 @@ As a result, the application must use the overload of `create_kernel_bundle_from_source` taking `std::vector` when creating a kernel bundle from this language. -=== Build options +=== Supported properties + +The `spirv` language supports only those properties from +link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[ +sycl_ext_oneapi_kernel_compiler] that are common to all source languages. This extension does not specify any options that may be passed via the -`build_options` property, however an implementation may allow -implementation-defined options to be passed this way. +`build_options` property to the `build` function, however an implementation may +allow implementation-defined options to be passed this way. === SPIR-V execution environment