26 inline namespace _V1 {
35 std::vector<device> DeviceList,
37 : MContext(Context), MDevices(DeviceList), MPropList(PropList) {
38 if (Context->getDevices().size() > 1) {
40 sycl::errc::feature_not_supported,
41 "multiple devices within a context are not supported with "
42 "sycl::program and sycl::kernel");
47 std::vector<std::shared_ptr<program_impl>> ProgramList,
50 MLinkOptions(LinkOptions), MBuildOptions(LinkOptions) {
52 if (ProgramList.empty()) {
53 throw runtime_error(
"Non-empty vector of programs expected",
54 PI_ERROR_INVALID_VALUE);
59 std::sort(ProgramList.begin(), ProgramList.end());
60 auto It = std::unique(ProgramList.begin(), ProgramList.end());
61 if (It != ProgramList.end()) {
62 throw runtime_error(
"Attempting to link a program with itself",
63 PI_ERROR_INVALID_PROGRAM);
66 MContext = ProgramList[0]->MContext;
67 if (MContext->getDevices().size() > 1) {
69 sycl::errc::feature_not_supported,
70 "multiple devices within a context are not supported with "
71 "sycl::program and sycl::kernel");
73 MDevices = ProgramList[0]->MDevices;
74 std::vector<device> DevicesSorted;
76 DevicesSorted = sort_devices_by_cl_device_id(MDevices);
78 check_device_feature_support<info::device::is_linker_available>(MDevices);
79 std::list<std::lock_guard<std::mutex>> Locks;
80 for (
const auto &Prg : ProgramList) {
81 Locks.emplace_back(Prg->MMutex);
83 if (Prg->MContext != MContext) {
84 throw invalid_object_error(
85 "Not all programs are associated with the same context",
86 PI_ERROR_INVALID_PROGRAM);
89 std::vector<device> PrgDevicesSorted =
90 sort_devices_by_cl_device_id(Prg->MDevices);
91 if (PrgDevicesSorted != DevicesSorted) {
92 throw invalid_object_error(
93 "Not all programs are associated with the same devices",
94 PI_ERROR_INVALID_PROGRAM);
100 std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
101 std::vector<sycl::detail::pi::PiProgram> Programs;
102 bool NonInterOpToLink =
false;
103 for (
const auto &Prg : ProgramList) {
104 if (!Prg->MLinkable && NonInterOpToLink)
106 NonInterOpToLink |= !Prg->MLinkable;
107 Programs.push_back(Prg->MProgram);
112 MContext->getHandleRef(), Devices.size(), Devices.data(),
113 LinkOptions.c_str(), Programs.size(), Programs.data(),
nullptr,
115 Plugin->checkPiResult<compile_program_error>(Err);
128 : MProgram(Program), MContext(Context), MLinkable(true) {
130 if (MProgram ==
nullptr) {
131 assert(InteropProgram &&
132 "No InteropProgram/PiProgram defined with piextProgramFromNative");
135 InteropProgram, MContext->getHandleRef(),
false, &MProgram);
144 std::vector<sycl::detail::pi::PiDevice> PiDevices(NumDevices);
148 PiDevices.data(),
nullptr);
150 std::vector<device> PlatformDevices =
151 MContext->getPlatformImpl()->get_devices();
155 auto NewEnd = std::remove_if(
156 PlatformDevices.begin(), PlatformDevices.end(),
158 return PiDevices.end() ==
159 std::find(PiDevices.begin(), PiDevices.end(),
160 detail::getSyclObjImpl(Dev)->getHandleRef());
162 PlatformDevices.erase(NewEnd, PlatformDevices.end());
163 MDevices = PlatformDevices;
164 assert(!MDevices.empty() &&
"No device found for this program");
170 sizeof(cl_program_binary_type), &BinaryType,
nullptr);
172 throw invalid_object_error(
173 "The native program passed to the program constructor has to be either "
174 "compiled or linked",
175 PI_ERROR_INVALID_PROGRAM);
180 std::vector<char> OptionsVector(Size);
183 OptionsVector.data(),
nullptr);
184 std::string Options(OptionsVector.begin(), OptionsVector.end());
185 switch (BinaryType) {
188 MCompileOptions = Options;
189 MBuildOptions = Options;
195 MBuildOptions = Options;
198 assert(
false &&
"BinaryType is invalid.");
211 if (!
is_host() && MProgram !=
nullptr) {
220 throw invalid_object_error(
221 "This instance of program doesn't support OpenCL interoperability.",
222 PI_ERROR_INVALID_PROGRAM);
225 return pi::cast<cl_program>(MProgram);
229 std::string CompileOptions) {
230 std::lock_guard<std::mutex> Lock(MMutex);
233 create_pi_program_with_kernel_name(
235 (!CompileOptions.empty()));
236 compile(CompileOptions);
242 std::lock_guard<std::mutex> Lock(MMutex);
245 check_device_feature_support<info::device::is_linker_available>(MDevices);
246 std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
250 LinkOpts = LinkOptions.c_str();
256 if (MProgram !=
nullptr)
261 MContext->getHandleRef(), Devices.size(), Devices.data(), LinkOpts,
262 1, &MProgram,
nullptr,
nullptr, &MProgram);
263 Plugin->checkPiResult<compile_program_error>(Err);
264 MLinkOptions = LinkOptions;
265 MBuildOptions = LinkOptions;
271 bool IsCreatedFromSource)
const {
274 return !IsCreatedFromSource;
277 std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
284 Device, MProgram, KernelName.c_str(), &function_ptr);
285 if (Err != PI_SUCCESS &&
286 Err != PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE &&
287 Err != PI_ERROR_INVALID_KERNEL_NAME)
289 "Error from piextGetDeviceFunctionPointer when called by program",
291 if (Err == PI_SUCCESS || Err == PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE)
299 std::shared_ptr<program_impl> PtrToSelf,
300 bool IsCreatedFromSource)
const {
303 if (IsCreatedFromSource)
304 throw invalid_object_error(
"This instance of program is a host instance",
305 PI_ERROR_INVALID_PROGRAM);
307 return createSyclObjFromImpl<kernel>(
308 std::make_shared<kernel_impl>(MContext, PtrToSelf));
310 auto [Kernel, ArgMask] = get_pi_kernel_arg_mask_pair(KernelName);
311 return createSyclObjFromImpl<kernel>(std::make_shared<kernel_impl>(
312 Kernel, MContext, PtrToSelf, IsCreatedFromSource,
nullptr, ArgMask));
320 std::vector<std::vector<char>> Result;
322 std::vector<size_t> BinarySizes(MDevices.size());
325 sizeof(size_t) * BinarySizes.size(), BinarySizes.data(),
nullptr);
327 std::vector<char *> Pointers;
328 for (
size_t I = 0; I < BinarySizes.size(); ++I) {
329 Result.emplace_back(BinarySizes[I]);
330 Pointers.push_back(Result[I].data());
333 sizeof(
char *) * Pointers.size(),
334 Pointers.data(),
nullptr);
338 void program_impl::compile(
const std::string &Options) {
339 check_device_feature_support<info::device::is_compiler_available>(MDevices);
340 std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
344 CompileOpts = Options.c_str();
348 MProgram, Devices.size(), Devices.data(), CompileOpts, 0,
nullptr,
349 nullptr,
nullptr,
nullptr);
351 if (Err != PI_SUCCESS) {
352 throw compile_program_error(
353 "Program compilation error:\n" +
357 MCompileOptions = Options;
358 MBuildOptions = Options;
361 void program_impl::build(
const std::string &Options) {
362 check_device_feature_support<info::device::is_compiler_available>(MDevices);
363 std::vector<sycl::detail::pi::PiDevice> Devices(get_pi_devices());
368 MProgram, Devices.size(), Devices.data(), Options.c_str(),
nullptr,
371 if (Err != PI_SUCCESS) {
372 throw compile_program_error(
373 "Program build error:\n" +
377 MBuildOptions = Options;
380 std::vector<sycl::detail::pi::PiDevice> program_impl::get_pi_devices()
const {
381 std::vector<sycl::detail::pi::PiDevice> PiDevices;
382 for (
const auto &Device : MDevices) {
388 std::pair<sycl::detail::pi::PiKernel, const KernelArgMask *>
389 program_impl::get_pi_kernel_arg_mask_pair(
const std::string &KernelName)
const {
390 std::pair<sycl::detail::pi::PiKernel, const KernelArgMask *> Result;
395 MProgram, KernelName.c_str(), &Result.first);
396 if (Err == PI_ERROR_INVALID_KERNEL_NAME) {
397 throw invalid_object_error(
398 "This instance of program does not contain the kernel requested",
401 Plugin->checkPiResult(Err);
413 program_impl::sort_devices_by_cl_device_id(std::vector<device> Devices) {
414 std::sort(Devices.begin(), Devices.end(),
415 [](
const device &id1,
const device &id2) {
416 return (detail::getSyclObjImpl(id1)->getHandleRef() <
417 detail::getSyclObjImpl(id2)->getHandleRef());
422 void program_impl::throw_if_state_is(
program_state State)
const {
423 if (MState == State) {
424 throw invalid_object_error(
"Invalid program state",
425 PI_ERROR_INVALID_PROGRAM);
429 void program_impl::throw_if_state_is_not(
program_state State)
const {
430 if (MState != State) {
431 throw invalid_object_error(
"Invalid program state",
432 PI_ERROR_INVALID_PROGRAM);
436 void program_impl::create_pi_program_with_kernel_name(
437 const std::string &KernelName,
bool JITCompilationIsRequired) {
438 assert(!MProgram &&
"This program already has an encapsulated PI program");
441 RTDeviceBinaryImage &Img = PM.getDeviceImage(
442 KernelName,
get_context(), FirstDevice, JITCompilationIsRequired);
443 MProgram = PM.createPIProgram(Img,
get_context(), {FirstDevice});
455 auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms();
458 for (SCItTy SCIt : SCRange) {
459 auto SCEntry = SpecConstRegistry.find((*SCIt)->Name);
460 if (SCEntry == SpecConstRegistry.end())
464 assert(SC.
isSet() &&
"uninitialized spec constant");
475 while (!Descriptors.
empty()) {
476 auto [Id, Offset, Size] =
477 Descriptors.
consume<uint32_t, uint32_t, uint32_t>();
The context class represents a SYCL context on which kernel functions may be executed.
void dropBytes(std::size_t Bytes)
ByteArray asByteArray() const
static ProgramManager & getInstance()
static std::string getProgramBuildLog(const sycl::detail::pi::PiProgram &Program, const ContextImplPtr Context)
void flushSpecConstants(const program_impl &Prg, pi::PiProgram NativePrg=nullptr, const RTDeviceBinaryImage *Img=nullptr)
Resolves given program to a device binary image and requests the program to flush constants the image...
const PropertyRange & getSpecConstants() const
Gets the iterator range over specialization constants in this binary image.
static const char * get()
void link(std::string LinkOptions="")
Links encapsulated raw program.
void compile_with_kernel_name(std::string KernelName, std::string CompileOptions)
Compiles the SYCL kernel function into the encapsulated raw program.
context get_context() const
ContextImplPtr getContextImplPtr() const
kernel get_kernel(std::string KernelName, std::shared_ptr< program_impl > PtrToSelf, bool IsCreatedFromSource) const
Returns a SYCL kernel for the SYCL kernel function defined by kernel name.
std::vector< std::vector< char > > get_binaries() const
Returns built program binaries.
std::vector< device > get_devices() const
void flush_spec_constants(const RTDeviceBinaryImage &Img, sycl::detail::pi::PiProgram NativePrg=nullptr) const
Takes current values of specialization constants and "injects" them into the underlying native progra...
const PluginPtr & getPlugin() const
pi_native_handle getNative() const
Returns the native plugin handle.
cl_program get() const
Returns a valid cl_program instance.
bool has_kernel(std::string KernelName, bool IsCreatedFromSource) const
Checks if kernel is available for this program.
sycl::detail::pi::PiProgram & getHandleRef()
const char * getValuePtr() const
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Provides an abstraction of a SYCL kernel.
Objects of the property_list class are containers for the SYCL properties.
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
std::shared_ptr< plugin > PluginPtr
std::lock_guard< SpinLock > LockGuard
pi_result piKernelCreate(pi_program program, const char *kernel_name, pi_kernel *ret_kernel)
uintptr_t pi_native_handle
pi_result piProgramGetBuildInfo(pi_program program, pi_device device, _pi_program_build_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle, pi_context context, bool pluginOwnsNativeHandle, pi_program *program)
Creates PI program object from a native handle.
pi_result piProgramRetain(pi_program program)
@ PI_USM_INDIRECT_ACCESS
indicates that the kernel might access data through USM ptrs
pi_result piProgramRelease(pi_program program)
pi_result piextProgramSetSpecializationConstant(pi_program prog, pi_uint32 spec_id, size_t spec_size, const void *spec_value)
Sets a specialization constant to a specific value.
pi_result piKernelSetExecInfo(pi_kernel kernel, pi_kernel_exec_info value_name, size_t param_value_size, const void *param_value)
API to set attributes controlling kernel execution.
pi_result piProgramGetInfo(pi_program program, pi_program_info param_name, size_t param_value_size, void *param_value, size_t *param_value_size_ret)
pi_result piProgramBuild(pi_program program, pi_uint32 num_devices, const pi_device *device_list, const char *options, void(*pfn_notify)(pi_program program, void *user_data), void *user_data)
pi_result piProgramLink(pi_context context, pi_uint32 num_devices, const pi_device *device_list, const char *options, pi_uint32 num_input_programs, const pi_program *input_programs, void(*pfn_notify)(pi_program program, void *user_data), void *user_data, pi_program *ret_program)
pi_result piProgramCompile(pi_program program, pi_uint32 num_devices, const pi_device *device_list, const char *options, pi_uint32 num_input_headers, const pi_program *input_headers, const char **header_include_names, void(*pfn_notify)(pi_program program, void *user_data), void *user_data)
pi_result piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle)
Gets the native handle of a PI program object.
@ PI_PROGRAM_INFO_NUM_DEVICES
@ PI_PROGRAM_INFO_BINARY_SIZES
@ PI_PROGRAM_INFO_BINARIES
@ PI_PROGRAM_INFO_DEVICES
pi_result piextGetDeviceFunctionPointer(pi_device device, pi_program program, const char *function_name, pi_uint64 *function_pointer_ret)
Retrieves a device function pointer to a user-defined function.
@ PI_PROGRAM_BINARY_TYPE_LIBRARY
@ PI_PROGRAM_BINARY_TYPE_EXECUTABLE
@ PI_PROGRAM_BINARY_TYPE_COMPILED_OBJECT
@ PI_PROGRAM_BINARY_TYPE_NONE
@ PI_PROGRAM_BUILD_INFO_OPTIONS
@ PI_PROGRAM_BUILD_INFO_BINARY_TYPE