Copyright © 2023 Intel Corporation. All rights reserved.
Khronos® is a registered trademark and SYCL™ and SPIR™ are trademarks of The Khronos Group Inc. OpenCL™ is a trademark of Apple Inc. used by permission by Khronos.
To report problems with this extension, please open a new issue at:
This extension is written against the SYCL 2020 revision 8 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision.
This extension also depends on the following other SYCL extensions:
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 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. This extension provides support for kernels written in SYCL according to the "free function kernel" syntax defined in 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.
Thus, an application can create a kernel bundle from a source string and then
build the bundle into "executable" bundle state.
Once the application obtains a kernel object, it can use existing APIs from
the core SYCL specification to set the value of kernel arguments and enqueue
the kernel to a device.
This extension provides a feature-test macro as described in the core SYCL
specification.
An implementation supporting this extension must predefine the macro
SYCL_EXT_ONEAPI_KERNEL_COMPILER
to one of the values defined in the table below.
Applications can test for the existence of this macro to determine if
the implementation supports this feature, or applications can test the macro’s
value to determine which of the extension’s features the implementation
supports.
| Value | Description |
|---|---|
1 |
The APIs of this experimental extension are not versioned, so the feature-test macro always has this value. |
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.
namespace sycl {
enum class bundle_state : /*unspecified*/ {
// ...
ext_oneapi_source
};
} // namespace syclThis extension adds the source_language enumeration, which identifies
possible languages for a kernel bundle that is in ext_oneapi_source state:
namespace sycl::ext::oneapi::experimental {
enum class source_language : /*unspecified*/ {
sycl
};
} // namespace sycl::ext::oneapi::experimentalThe 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
sycl_ext_oneapi_free_function_kernels.
Other extensions may provide other enumerators that correspond to other
languages.
This extension adds the following new member functions to the device class:
Returns: The value |
|
Returns: The value |
This extension adds the following new free functions to create and build a
kernel bundle in ext_oneapi_source state.
Constraints: Available only when Effects: Creates a new kernel bundle that represents a kernel written in the
source language Each source language Returns: The newly created kernel bundle, which has Throws:
[Note: Calling this function does not attempt to compile the source code.
As a result, syntax errors in This function succeeds even if some devices in |
|
Constraints: Available only when Effects (1): The source code from The new bundle has the same associated context as Effects (2): Equivalent to
Returns: The newly created kernel bundle, which has Throws:
[Note: An uncaught |
|
Constraints: Available only when Effects (1): The source code from The new bundle has the same associated context as Effects (2): Equivalent to
Returns: The newly created kernel bundle, which has Throws:
[Note: An uncaught |
This extension adds the following properties, which can be used in conjunction
with the create_kernel_bundle_from_source function that is defined above:
This property provides the name and content of include files that can be
referenced from the source code in the When the source language is [Note: This property is only required if an Effects (1): Creates a new Effects (2): Creates a new Effects (3): Adds a (Name, Content) pair to the property. Throws (3):
|
This extension adds the following properties, which can be used in conjunction
with the build and compile function that is defined above:
This property provides build options that may affect the compilation or linking
of the kernel, where each build option is a string.
All source languages support the Effects (1): Constructs a Effects (2): Constructs a Effects (3): Constructs a Effects (4): Adds |
|
This property allows the caller to request a log to be created with additional
information about the compilation and linking operations.
Use of this property is not required in order to get information about a failed
build.
When a build fails, an Instead, the Effects (1): Constructs a Remarks (1): When |
|
This property is useful when the source language represents names differently in
the source code and the generated code.
For example, C++ function names and the names of static variables at global
scope are "mangled" in an implementation-defined way in the generated code.
The precise meaning of this property is defined by each source language, but in
general it allows the application to supply a list of 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 Effects (1): Creates a new Effects (2): Creates a new Effects (3): Creates a new Effects (4): Adds Preconditions (2-4): Each source language defines its own requirements for the
registered names.
For the language [Note: It is not an error to have duplicate names in a |
This extension adds the following new free functions to link kernel bundles in
object state. This differs from the regular sycl::link function the
additional property list argument.
Constraints: Available only when Effects: Duplicate device images from The new bundle has the same associated context as those in Returns: The new kernel bundle. Throws:
|
|
Effects (1): Equivalent to Effects (2): Equivalent to Effects (3): Equivalent to
|
This extension adds the following properties, which can be used in conjunction
with the link function that is defined above:
This property instructs the Fast linking offers potentially faster link times, at the potential cost of slower kernel execution time. The latter effect would mainly be due to link-time optimizations the device compiler could have done during regular linking, that cannot be done when doing fast linking. For example, the binaries produced may contain ahead-of-time compiled binary images together with just-in-time compiled binary images, with the kernels and exported functions potentially overlapping. When fast-linking is enabled, the implementation will try to use the ahead-of-time compiled binary images over their just-in-time compiled counterparts. Effects (1): Creates a new |
This extension adds the following constraint to some of the kernel_bundle
member functions from the core SYCL specification:
Constraints: This function is not available when
Stateisbundle_state::ext_oneapi_source.
This new constraint applies to the following member functions:
-
empty; -
All overloads and function templates of
has_kernel; -
get_kernel_ids; -
contains_specialization_constants; -
native_specialization_constant; -
has_specialization_constant; -
get_specialization_constant; -
begin; and -
end.
As a result, the only kernel_bundle member functions from the core SYCL
specification that are available for bundles in ext_oneapi_source state are
get_backend, get_context, and get_devices.
This extension adds the following constraint to the join functions from the
core SYCL specification:
Constraints: This function is not available when
Stateisbundle_state::ext_oneapi_source.
Kernels created from online compilation of source code do not have any
associated kernel_id.
Therefore, the function kernel_bundle::get_kernel_ids returns an empty vector
of kernel_id objects if the kernel bundle was created from a bundle of state
bundle_state::ext_oneapi_source.
This extensions adds the following new kernel_bundle member functions:
namespace sycl {
template <bundle_state State>
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
Constraints: This function is not available when Returns: The value |
|
Constraints: This function is available only when Returns: A Throws:
|
|
Constraints: This function is not available when Returns: If the kernel bundle was created from a bundle of state
Throws:
|
When the source language is source_language::sycl, the compiler searches
multiple locations to find files referenced by #include statements.
Any include files defined via the include_files property are searched first,
followed by the directories below, in order:
-
The current working directory.
-
Any directory added explicitly to the search list via the
build_optionsproperty.
Finally, the compiler searches a set of implicitly available header files, which
do not need to be specified via the include_files property:
-
<sycl/sycl.hpp>; -
The C++ standard library headers;
-
The SYCL backend headers
"sycl/backend/<backend_name>.hpp"for any backends that the implementation supports; and -
Any SYCL extension headers in
"sycl/ext"for extensions that the implementation supports.
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 C++ mangled name)
or by using the registered_names property.
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:
std::string source = R"""(
#include <sycl/sycl.hpp>
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:
sycl::kernel_bundle<sycl::bundle_state::executable> kb = /*...*/;
sycl::kernel k = kb.ext_oneapi_get_kernel("foo");When the kernel is not declared as extern "C", the compiler generates a
mangled name, so it is more convenient to use the registered_names property.
Each string in the property must be the C++ 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:
std::string source = R"""(
#include <sycl/sycl.hpp>
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:
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src = /*...*/;
sycl::kernel_bundle<sycl::bundle_state::executable> kb = syclexp::build(kb_src,
syclexp::properties{syclexp::registered_names{"mykernels::bar"}});
sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar");The C++ 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:
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src = /*...*/;
sycl::kernel_bundle<sycl::bundle_state::executable> kb = syclexp::build(kb_src,
syclexp::properties{syclexp::registered_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_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.
The registered_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:
std::string source = R"""(
#include <sycl/sycl.hpp>
namespace syclexp = sycl::ext::oneapi::experimental;
template<typename T>
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
void bartmpl(T *in, T *out) {/*...*/}
)""";The application can use the registered_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:
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src = /*...*/;
sycl::kernel_bundle<sycl::bundle_state::executable> kb = syclexp::build(kb_src,
syclexp::properties{syclexp::registered_names{{"bartmpl<float>", "bartmpl<int>"}});
sycl::kernel k_float = kb.ext_oneapi_get_kernel("bartmpl<float>");
sycl::kernel k_int = kb.ext_oneapi_get_kernel("bartmpl<int>");The following example demonstrates how a SYCL application can define a kernel as a string and then compile and launch it.
#include <sycl/sycl.hpp>
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 a kernel, defined as a SYCL "free function kernel".
std::string source = R"""(
#include <sycl/sycl.hpp>
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<1>().get_global_linear_id();
ptr[id] = start + static_cast<float>(id);
}
)""";
// Create a kernel bundle in "source" state.
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> kb_src =
syclexp::create_kernel_bundle_from_source(
q.get_context(),
syclexp::source_language::sycl,
source);
// Compile the kernel. There is no need to use the "registered_names"
// property because the kernel is declared extern "C".
sycl::kernel_bundle<sycl::bundle_state::executable> kb_exe =
syclexp::build(kb_src);
// Get the kernel via its compiler-generated name.
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel("iota");
float *ptr = sycl::malloc_shared<float>(NUM, q);
q.submit([&](sycl::handler &cgh) {
// Set the values of the kernel arguments.
cgh.set_args(3.14f, 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();
sycl::free(ptr, q);
}This example demonstrates how to use the registered_names property to
disambiguate a kernel function that has several overloads.
#include <sycl/sycl.hpp>
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 <sycl/sycl.hpp>
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<1>().get_global_linear_id();
ptr[id] = start + static_cast<float>(id);
}
SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
void iota(int start, int *ptr) {
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
ptr[id] = start + static_cast<int>(id);
}
)""";
// Create a kernel bundle in "source" state.
sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source> 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<sycl::bundle_state::executable> kb_exe =
syclexp::build(kb_src, syclexp::properties{syclexp::registered_names{iota_name}});
// Get the kernel by passing the same string we used to construct the
// "registered_names" property.
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name);
int *ptr = sycl::malloc_shared<int>(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();
sycl::free(ptr, q);
}-
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 usetypeidto get astd::type_info. You can then calltype_info::nameto get an implementation-defined name for the type. For clang on Linux, this returns the type’s mangled name. You can then callabi::__cxa_demangleto get an unmangled name for the type. I’m not sure about the details on Windows hosts, though. Iftype_info::namereturns a mangled name on Windows too, then maybe we can still useabi::__cxa_demangleto 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 DPC++. 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.
The SYCL runtime compiler supports the following DPC++ options to be passed in
the build_options property.
Some options have equivalent long (starting with --) and short (starting with
-) option names. When using the long option name, an argument can be either
separated by = in the same element of the build_options property, or given
as a separate element. When using the short option name, an argument is either
appended directly after the option name, or given as a separate element in the
build_options property. The following example shows how to construct the
build_options property with each of the forms.
build_options{{
{"--include-directory=dir1"},
{"--include-directory"}, {"dir2"},
{"-Idir3"},
{"-I"}, {"dir4"}
}};Add <dir> to to the search list for include files (see section "Including
files when the language is sycl"). This is useful, for example, to compile
kernels using external libraries.
Define macro <name>, optionally to the given <value>.
The build_options property accepts warning (-W) and remark (-R) emission
options supported by the clang compiler. For an overview of these options, see
https://clang.llvm.org/docs/DiagnosticsReference.html. The specific options
available for SYCL runtime compilation depend on the version of the DPC++
compiler distributed with the SYCL runtime used by the application.
Note: Use the save_log property to obtain detailed output from the compilation
process.
Pass <arg> to the backend of the device compiler. When using -Xs<arg>, a -
is prepended to <arg> before handing it to the backend. Otherwise, <arg> is
passed on unmodified.
For example, the following forms are equivalent:
build_options{{
{"-XsDFOO=bar"},
{"-Xs"}, {"-DFOO=bar"}
}};Relax the requirement that parameter types for free-function kernels must be forward-declarable.
The first time this option is passed, the compiler finds the initial set of
preprocessor directives (e.g., #define/#include) and comments in the
compiled source string (the preamble) and precompiles it. Essentialy, it behaves
like a precompiled header containing that preamble. On subsequent compilations,
if the compiled source string has the same preamble and the same compilation
options are used, the precompiled preamble is used, which speeds up compilation.
If the compiled source string has a different preamble or compilation options differ, a new precompiled preamble is generated, and that preamble can also be used to speed up subsequent compilations. These precompiled preambles are stored internally in memory, so they do not persist from one execution of the application to the next.
The preamble ends with the first statement that is not a preprocessor directive or a comment. For example, in the code below, the preamble ends immediately before the namespace syclext = statement.
#define SYCL_SIMPLE_SWIZZLES
#include <sycl/sycl.hpp>
// Auto-detected preamble ends before next line:
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(sycl::vec<int, 2> *p) {
size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id();
p[id] = p[id].xx();
}The compiler uses the following factors when deciding whether a previously generated precompiled preamble can be used:
-
The preamble must exactly match (including whitespace and comments).
-
The compilation options must match (including the same order and the same spelling).
-
There are also certain restrictions that the user must avoid:
-
The content of each header file in the preamble must not change from one compilation to another.
-
It is not recommended to use the
__DATE__or__TIME__macros in the preamble header files. Depending on the circumstances, these macros may be replaced with the date / time that corresponds to the time at which the precompiled preamble was generated, rather than the time at which the source string is compiled. See also the clang compiler options-Wpch-date-timeand-Werror=pch-date-time, which cause the compiler to diagnose a warning or error in this scenario.
-
Similar to --auto-pch but the precompiled preamble is stored on the filesystem
at the location <dir>. That allows an application to precompile the preamble
once and reuse it across multiple invocation of the application which isn’t
possible with --auto-pch that stores the precompiled preamble in-memory. This
option is incompatible with --auto-pch.
Some notes about the current behavior:
-
Threads only communicate via the persistent cache on the filesystem with no other inter-thread communication to avoid extra synchronization overhead. In an unlikely event that a preamble is missing from the cache and multiple threads need it, each thread will build the preamble.
-
In the very unlikely event of a hash collision the precompiled preamble from the persistent cache is ignored and the source is compiled without it.
-
Unlike
--auto-pch, the--persistent-auto-pchoption indexes the precompiled preamble by a hash of the preamble content/compilation options, as opposed to using those directly as a key. As a result, there is a highly unlikely possibility that two different preambles will produce the same hash value. Even if this occurs, the compilation won’t fail. However, a compilation using the conflicting hash would proceed without pre-compiled preamble support as if this option wasn’t enabled.
As the DPC++ frontend is integrated tightly in the runtime compilation
pipeline, the application cannot change the runtime compiler’s action (e.g.
-c, -S) or output file (-o). Similarly, options related to linking (e.g.
-L) are incompatible, including the SYCL-specific -fsycl-link action. The
implementation throws an exception with the errc::invalid error code when it
detects an option that conflicts with the runtime compilation pipeline.
The kernels in a SYCL source string are compiled automatically to native code
for all devices passed to the build function (see section "New free functions
to create and build kernel bundles"). The implementation rejects the use of the
-fsycl-targets= options to request ahead-of-time (AOT) compilation, and throws
an exception with the errc::invalid error code when this option is detected.
The application can use the -Xs option described above to pass options to the
backend of the device compiler, but all other options to control AOT compilation
are ignored.
The SYCL runtime compiler currently does not implement the logic required to
support the -fno-sycl-device-code-split-esimd option, and throws an
exception with the errc::invalid error code when this option is detected. As
a consequence, the invoke_simd functionality is unavailable. However, the SYCL
runtime compiler supports ESIMD kernels and source strings containing a mix of
SYCL and ESIMD kernels.
The implementation currently lacks the necessary linking of device libraries to
support device, memory and thread sanitizers for runtime-compiled code. If the
-fsanitize= option is detected, an exception with the errc::invalid error
code is thrown. Other means of activating the sanitizer (e.g. via
-Xsycl-device-frontend) may cause the runtime compilation to fail.
The kernel_compiler implementation in DPC++ supports persistent caching. To
enable it, set the the environment variable SYCL_CACHE_PERSISTENT=1. The
location of the cache can be changed by setting SYCL_CACHE_DIR. Refer to
https://intel.github.io/llvm/design/KernelProgramCache.html#persistent-cache for
more details on how to control the cache.