20 inline namespace _V1 {
23 : impl(
std::make_shared<detail::kernel_id_impl>(Name)) {}
25 const char *kernel_id::get_name() const
noexcept {
return impl->get_name(); }
34 return impl->has_kernel(KernelID);
37 bool device_image_plain::has_kernel(
const kernel_id &KernelID,
39 return impl->has_kernel(KernelID, Dev);
43 return impl->getNative();
50 bool kernel_bundle_plain::empty() const
noexcept {
return impl->empty(); }
53 return impl->get_backend();
57 return impl->get_context();
60 std::vector<device> kernel_bundle_plain::get_devices() const
noexcept {
61 return impl->get_devices();
65 return impl->get_kernel_ids();
68 bool kernel_bundle_plain::contains_specialization_constants() const
noexcept {
69 return impl->contains_specialization_constants();
72 bool kernel_bundle_plain::native_specialization_constant() const
noexcept {
73 return impl->native_specialization_constant();
77 return impl->get_kernel(KernelID, impl);
89 return impl->has_kernel(KernelID);
92 bool kernel_bundle_plain::has_kernel(
const kernel_id &KernelID,
94 return impl->has_kernel(KernelID, Dev);
97 bool kernel_bundle_plain::has_specialization_constant_impl(
98 const char *SpecName)
const noexcept {
99 return impl->has_specialization_constant(SpecName);
102 void kernel_bundle_plain::set_specialization_constant_impl(
103 const char *SpecName,
void *Value,
size_t Size)
noexcept {
104 impl->set_specialization_constant_raw_value(SpecName, Value, Size);
107 void kernel_bundle_plain::get_specialization_constant_impl(
108 const char *SpecName,
void *Value)
const noexcept {
109 impl->get_specialization_constant_raw_value(SpecName, Value);
112 bool kernel_bundle_plain::is_specialization_constant_set(
113 const char *SpecName)
const noexcept {
114 return impl->is_specialization_constant_set(SpecName);
117 bool kernel_bundle_plain::ext_oneapi_has_kernel(
const std::string &name) {
118 return impl->ext_oneapi_has_kernel(name);
122 return impl->ext_oneapi_get_kernel(name, impl);
129 const std::vector<device>
131 std::vector<device> UniqueDevices;
134 std::unordered_set<device> UniqueDeviceSet;
135 for (
const device &Dev : Devs)
136 if (UniqueDeviceSet.insert(Dev).second)
137 UniqueDevices.push_back(Dev);
139 return UniqueDevices;
143 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
144 return detail::ProgramManager::getInstance().getSYCLKernelID(
147 return detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
154 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, State);
159 const std::vector<kernel_id> &KernelIDs,
161 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, KernelIDs,
168 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs, Selector,
174 const std::vector<device> &Devs) {
175 return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
178 std::shared_ptr<detail::kernel_bundle_impl>
179 join_impl(
const std::vector<detail::KernelBundleImplPtr> &Bundles,
181 return std::make_shared<detail::kernel_bundle_impl>(Bundles, State);
188 if (Devs.empty() || !AllDevicesInTheContext)
190 "Not all devices are associated with the context or "
191 "vector of devices is empty");
193 if (bundle_state::input == State &&
196 if (bundle_state::object == State &&
200 const std::vector<device_image_plain> DeviceImages =
201 detail::ProgramManager::getInstance()
202 .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
204 return (
bool)DeviceImages.size();
208 const std::vector<kernel_id> &KernelIds,
213 if (Devs.empty() || !AllDevicesInTheContext)
215 "Not all devices are associated with the context or "
216 "vector of devices is empty");
218 bool DeviceHasRequireAspectForState =
true;
219 if (bundle_state::input == State) {
220 DeviceHasRequireAspectForState =
222 return Dev.has(aspect::online_compiler);
224 }
else if (bundle_state::object == State) {
225 DeviceHasRequireAspectForState =
227 return Dev.has(aspect::online_linker);
231 if (!DeviceHasRequireAspectForState)
234 const std::vector<device_image_plain> DeviceImages =
235 detail::ProgramManager::getInstance()
236 .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
238 std::set<kernel_id, LessByNameComp> CombinedKernelIDs;
240 const std::shared_ptr<device_image_impl> &DeviceImageImpl =
243 CombinedKernelIDs.insert(DeviceImageImpl->get_kernel_ids_ptr()->begin(),
244 DeviceImageImpl->get_kernel_ids_ptr()->end());
247 const bool AllKernelIDsRepresented =
249 [&CombinedKernelIDs](
const kernel_id &KernelID) {
250 return CombinedKernelIDs.count(KernelID);
253 return AllKernelIDsRepresented;
256 std::shared_ptr<detail::kernel_bundle_impl>
258 const std::vector<device> &Devs,
const property_list &PropList) {
259 return std::make_shared<detail::kernel_bundle_impl>(
260 InputBundle, Devs, PropList, bundle_state::object);
263 std::shared_ptr<detail::kernel_bundle_impl>
265 const std::vector<device> &Devs,
const property_list &PropList) {
266 return std::make_shared<detail::kernel_bundle_impl>(ObjectBundles, Devs,
270 std::shared_ptr<detail::kernel_bundle_impl>
272 const std::vector<device> &Devs,
const property_list &PropList) {
273 return std::make_shared<detail::kernel_bundle_impl>(
274 InputBundle, Devs, PropList, bundle_state::executable);
281 std::vector<sycl::device> IntersectDevices;
282 std::vector<unsigned int> DevsCounters;
283 std::map<device, unsigned int, LessByHash<device>> DevCounters;
288 DevCounters[Device]++;
292 for (
const std::pair<const device, unsigned int> &It : DevCounters)
293 if (ObjectBundles.size() == It.second)
294 IntersectDevices.push_back(It.first);
296 return IntersectDevices;
306 return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
310 if (KernelIDs.empty())
315 auto doesImageTargetMatchDevice = [](
const device &Dev,
316 const detail::RTDeviceBinaryImage &Img) {
317 const char *Target = Img.getRawData().DeviceTargetSpec;
320 if (BE == sycl::backend::ext_intel_esimd_emulator) {
322 return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0));
325 return (BE == sycl::backend::opencl ||
326 BE == sycl::backend::ext_oneapi_level_zero);
332 return Dev.
is_gpu() && (BE == sycl::backend::opencl ||
333 BE == sycl::backend::ext_oneapi_level_zero);
338 return BE == sycl::backend::ext_oneapi_cuda;
340 return BE == sycl::backend::ext_oneapi_hip;
350 for (
const auto &KernelID : KernelIDs) {
351 std::set<detail::RTDeviceBinaryImage *> BinImages =
352 detail::ProgramManager::getInstance().getRawDeviceImages({KernelID});
355 [&](
const detail::RTDeviceBinaryImage *Img) {
356 return doesDevSupportDeviceRequirements(Dev, *Img) &&
357 doesImageTargetMatchDevice(Dev, *Img);
368 namespace ext::oneapi::experimental {
379 bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) ||
380 (BE == sycl::backend::opencl);
383 if (Language == source_language::opencl) {
385 }
else if (Language == source_language::spirv) {
407 "kernel_bundle creation from source not supported");
409 std::shared_ptr<kernel_bundle_impl> KBImpl =
410 std::make_shared<kernel_bundle_impl>(SyclContext, Language, Source);
411 return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
417 const std::vector<std::byte> &Bytes) {
421 "kernel_bundle creation from source not supported");
423 std::shared_ptr<kernel_bundle_impl> KBImpl =
424 std::make_shared<kernel_bundle_impl>(SyclContext, Language, Bytes);
425 return sycl::detail::createSyclObjFromImpl<source_kb>(KBImpl);
434 const std::vector<device> &Devices,
437 std::vector<device> UniqueDevices =
439 std::shared_ptr<kernel_bundle_impl> sourceImpl =
getSyclObjImpl(SourceKB);
440 std::shared_ptr<kernel_bundle_impl> KBImpl =
441 sourceImpl->build_from_source(UniqueDevices,
BuildOptions, LogPtr);
442 return sycl::detail::createSyclObjFromImpl<exe_kb>(KBImpl);
The context class represents a SYCL context on which kernel functions may be executed.
backend get_backend() const noexcept
Returns the backend associated with this context.
const char * data() const noexcept
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
bool is_accelerator() const
Check if device is an accelerator device.
bool is_gpu() const
Check if device is a GPU device.
static std::vector< device > get_devices(info::device_type deviceType=info::device_type::all)
Query available SYCL devices.
backend get_backend() const noexcept
Returns the backend associated with this device.
bool is_cpu() const
Check if device is a CPU device.
The kernel_bundle class represents collection of device images in a particular state.
Objects of the class identify kernel is some kernel_bundle related APIs.
Provides an abstraction of a SYCL kernel.
Objects of the property_list class are containers for the SYCL properties.
std::vector< sycl::device > find_device_intersection(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles)
std::shared_ptr< detail::kernel_bundle_impl > build_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
static bool checkAllDevicesAreInContext(const std::vector< device > &Devices, const context &Context)
std::function< bool(const detail::DeviceImageImplPtr &DevImgImpl)> DevImgSelectorImpl
static bool checkAllDevicesHaveAspect(const std::vector< device > &Devices, aspect Aspect)
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
const std::vector< device > removeDuplicateDevices(const std::vector< device > &Devs)
detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
std::shared_ptr< detail::kernel_bundle_impl > compile_impl(const kernel_bundle< bundle_state::input > &InputBundle, const std::vector< device > &Devs, const property_list &PropList)
std::shared_ptr< detail::kernel_bundle_impl > join_impl(const std::vector< detail::KernelBundleImplPtr > &Bundles, bundle_state State)
kernel_id get_kernel_id_impl(string_view KernelName)
std::shared_ptr< detail::kernel_bundle_impl > KernelBundleImplPtr
bool has_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs, bundle_state State)
detail::KernelBundleImplPtr get_empty_interop_kernel_bundle_impl(const context &Ctx, const std::vector< device > &Devs)
std::shared_ptr< detail::kernel_bundle_impl > link_impl(const std::vector< kernel_bundle< bundle_state::object >> &ObjectBundles, const std::vector< device > &Devs, const property_list &PropList)
bool OpenCLC_Compilation_Available()
exe_kb build_from_source(source_kb &SourceKB, const std::vector< device > &Devices, const std::vector< std::string > &BuildOptions, std::string *LogPtr)
bool is_source_kernel_bundle_supported(backend BE, source_language Language)
kernel_bundle< bundle_state::ext_oneapi_source > create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::string &Source)
sycl::detail::kernel_bundle_impl kernel_bundle_impl
kernel_bundle< sycl::bundle_state::ext_oneapi_source > source_kb
std::vector< kernel_id > get_kernel_ids()
bool is_compatible(const std::vector< kernel_id > &KernelIDs, const device &Dev)
std::error_code make_error_code(sycl::errc E) noexcept
Constructs an error code using e and sycl_category()
uintptr_t pi_native_handle
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64
Device-specific binary images produced from SPIR-V 64-bit <-> various "spir64_*" triples for specific...
#define __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64
PTX 64-bit image <-> "nvptx64", 64-bit NVIDIA PTX device.
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64
SPIR-V 64-bit image <-> "spir64", 64-bit OpenCL device.
#define __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA
bool all_of(const simd_mask< _Tp, _Abi > &) noexcept
_Abi const simd< _Tp, _Abi > & noexcept
bool none_of(const simd_mask< _Tp, _Abi > &) noexcept