diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index e72cdebf506c2..70b2e80c35dd0 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -223,7 +223,6 @@ set(SYCL_COMMON_SOURCES "detail/memory_manager.cpp" "detail/pipes.cpp" "detail/platform_impl.cpp" - "detail/program_impl.cpp" "detail/program_manager/program_manager.cpp" "detail/queue_impl.cpp" "detail/online_compiler/online_compiler.cpp" diff --git a/sycl/source/backend/level_zero.cpp b/sycl/source/backend/level_zero.cpp index ade630ba426c2..19ce9281f289d 100644 --- a/sycl/source/backend/level_zero.cpp +++ b/sycl/source/backend/level_zero.cpp @@ -8,7 +8,6 @@ #include #include -#include #include #include #include diff --git a/sycl/source/backend/opencl.cpp b/sycl/source/backend/opencl.cpp index 9bdd8bfaeb317..07aec7ba0549c 100644 --- a/sycl/source/backend/opencl.cpp +++ b/sycl/source/backend/opencl.cpp @@ -9,7 +9,6 @@ #include #include #include -#include #include #include diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp index 8502f3489b9c7..545e4701f360c 100644 --- a/sycl/source/detail/kernel_impl.cpp +++ b/sycl/source/detail/kernel_impl.cpp @@ -9,7 +9,6 @@ #include #include #include -#include #include @@ -18,34 +17,14 @@ inline namespace _V1 { namespace detail { kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, - ContextImplPtr Context, - KernelBundleImplPtr KernelBundleImpl, - const KernelArgMask *ArgMask) - : kernel_impl(Kernel, Context, - std::make_shared(Context, Kernel), - /*IsCreatedFromSource*/ true, KernelBundleImpl, ArgMask) { - // Enable USM indirect access for interoperability kernels. - // Some PI Plugins (like OpenCL) require this call to enable USM - // For others, PI will turn this into a NOP. - if (Context->getPlatformImpl()->supports_usm()) - getPlugin()->call( - MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE); - - // This constructor is only called in the interoperability kernel constructor. - MIsInterop = true; -} - -kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, - ContextImplPtr ContextImpl, ProgramImplPtr ProgramImpl, - bool IsCreatedFromSource, + ContextImplPtr ContextImpl, KernelBundleImplPtr KernelBundleImpl, const KernelArgMask *ArgMask) : MKernel(Kernel), MContext(ContextImpl), - MProgram(ProgramImpl->getHandleRef()), - MCreatedFromSource(IsCreatedFromSource), - MKernelBundleImpl(std::move(KernelBundleImpl)), - MKernelArgMaskPtr{ArgMask} { - + MProgram(ProgramManager::getInstance().getPiProgramFromPiKernel( + Kernel, ContextImpl)), + MCreatedFromSource(true), MKernelBundleImpl(std::move(KernelBundleImpl)), + MIsInterop(true), MKernelArgMaskPtr{ArgMask} { sycl::detail::pi::PiContext Context = nullptr; // Using the plugin from the passed ContextImpl getPlugin()->call( @@ -55,7 +34,12 @@ kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, "Input context must be the same as the context of cl_kernel", PI_ERROR_INVALID_CONTEXT); - MIsInterop = ProgramImpl->isInterop(); + // Enable USM indirect access for interoperability kernels. + // Some PI Plugins (like OpenCL) require this call to enable USM + // For others, PI will turn this into a NOP. + if (ContextImpl->getPlatformImpl()->supports_usm()) + getPlugin()->call( + MKernel, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE); } kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, @@ -71,9 +55,6 @@ kernel_impl::kernel_impl(sycl::detail::pi::PiKernel Kernel, MIsInterop = MKernelBundleImpl->isInterop(); } -kernel_impl::kernel_impl(ContextImplPtr Context, ProgramImplPtr ProgramImpl) - : MContext(Context), MProgram(ProgramImpl->getHandleRef()) {} - kernel_impl::~kernel_impl() { try { // TODO catch an exception and put it to list of asynchronous exceptions diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp index 1a1542d0d409b..29f716806f8cb 100644 --- a/sycl/source/detail/kernel_impl.hpp +++ b/sycl/source/detail/kernel_impl.hpp @@ -26,11 +26,9 @@ namespace sycl { inline namespace _V1 { namespace detail { // Forward declaration -class program_impl; class kernel_bundle_impl; using ContextImplPtr = std::shared_ptr; -using ProgramImplPtr = std::shared_ptr; using KernelBundleImplPtr = std::shared_ptr; using sycl::detail::pi::PiProgram; class kernel_impl { @@ -38,8 +36,7 @@ class kernel_impl { /// Constructs a SYCL kernel instance from a PiKernel /// /// This constructor is used for plug-in interoperability. It always marks - /// kernel as being created from source and creates a new program_impl - /// instance. + /// kernel as being created from source. /// /// \param Kernel is a valid PiKernel instance /// \param Context is a valid SYCL context @@ -48,24 +45,6 @@ class kernel_impl { KernelBundleImplPtr KernelBundleImpl, const KernelArgMask *ArgMask = nullptr); - /// Constructs a SYCL kernel instance from a SYCL program and a PiKernel - /// - /// This constructor creates a new instance from PiKernel and saves - /// the provided SYCL program. If context of PiKernel differs from - /// context of the SYCL program, an invalid_parameter_error exception is - /// thrown. - /// - /// \param Kernel is a valid PiKernel instance - /// \param ContextImpl is a valid SYCL context - /// \param ProgramImpl is a valid instance of program_impl - /// \param IsCreatedFromSource is a flag that indicates whether program - /// is created from source code - /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl - kernel_impl(sycl::detail::pi::PiKernel Kernel, ContextImplPtr ContextImpl, - ProgramImplPtr ProgramImpl, bool IsCreatedFromSource, - KernelBundleImplPtr KernelBundleImpl, - const KernelArgMask *ArgMask); - /// Constructs a SYCL kernel_impl instance from a SYCL device_image, /// kernel_bundle and / PiKernel. /// @@ -78,12 +57,6 @@ class kernel_impl { const KernelArgMask *ArgMask, PiProgram ProgramPI, std::mutex *CacheMutex); - /// Constructs a SYCL kernel for host device - /// - /// \param Context is a valid SYCL context - /// \param ProgramImpl is a valid instance of program_impl - kernel_impl(ContextImplPtr Context, ProgramImplPtr ProgramImpl); - // This section means the object is non-movable and non-copyable // There is no need of move and copy constructors in kernel_impl. // If they need to be added, piKernelRetain method for MKernel diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp deleted file mode 100644 index ed78ecb440ef7..0000000000000 --- a/sycl/source/detail/program_impl.cpp +++ /dev/null @@ -1,473 +0,0 @@ -//==----- program_impl.cpp --- SYCL program implementation -----------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -namespace sycl { -inline namespace _V1 { -namespace detail { - -program_impl::program_impl(ContextImplPtr Context, - const property_list &PropList) - : program_impl(Context, Context->get_info(), - PropList) {} - -program_impl::program_impl(ContextImplPtr Context, - std::vector DeviceList, - const property_list &PropList) - : MContext(Context), MDevices(DeviceList), MPropList(PropList) { - if (Context->getDevices().size() > 1) { - throw sycl::exception( - sycl::errc::feature_not_supported, - "multiple devices within a context are not supported with " - "sycl::program and sycl::kernel"); - } -} - -program_impl::program_impl( - std::vector> ProgramList, - std::string LinkOptions, const property_list &PropList) - : MState(program_state::linked), MPropList(PropList), - MLinkOptions(LinkOptions), MBuildOptions(LinkOptions) { - // Verify arguments - if (ProgramList.empty()) { - throw runtime_error("Non-empty vector of programs expected", - PI_ERROR_INVALID_VALUE); - } - - // Sort the programs to avoid deadlocks due to locking multiple mutexes & - // verify that all programs are unique. - std::sort(ProgramList.begin(), ProgramList.end()); - auto It = std::unique(ProgramList.begin(), ProgramList.end()); - if (It != ProgramList.end()) { - throw runtime_error("Attempting to link a program with itself", - PI_ERROR_INVALID_PROGRAM); - } - - MContext = ProgramList[0]->MContext; - if (MContext->getDevices().size() > 1) { - throw sycl::exception( - sycl::errc::feature_not_supported, - "multiple devices within a context are not supported with " - "sycl::program and sycl::kernel"); - } - MDevices = ProgramList[0]->MDevices; - std::vector DevicesSorted; - DevicesSorted = sort_devices_by_cl_device_id(MDevices); - - check_device_feature_support(MDevices); - std::list> Locks; - for (const auto &Prg : ProgramList) { - Locks.emplace_back(Prg->MMutex); - Prg->throw_if_state_is_not(program_state::compiled); - if (Prg->MContext != MContext) { - throw invalid_object_error( - "Not all programs are associated with the same context", - PI_ERROR_INVALID_PROGRAM); - } - - std::vector PrgDevicesSorted = - sort_devices_by_cl_device_id(Prg->MDevices); - if (PrgDevicesSorted != DevicesSorted) { - throw invalid_object_error( - "Not all programs are associated with the same devices", - PI_ERROR_INVALID_PROGRAM); - } - } - - std::vector Devices(get_pi_devices()); - std::vector Programs; - bool NonInterOpToLink = false; - for (const auto &Prg : ProgramList) { - if (!Prg->MLinkable && NonInterOpToLink) - continue; - NonInterOpToLink |= !Prg->MLinkable; - Programs.push_back(Prg->MProgram); - } - const PluginPtr &Plugin = getPlugin(); - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MContext->getHandleRef(), Devices.size(), Devices.data(), - LinkOptions.c_str(), Programs.size(), Programs.data(), nullptr, - nullptr, &MProgram); - Plugin->checkPiResult(Err); -} - -program_impl::program_impl(ContextImplPtr Context, - pi_native_handle InteropProgram) - : program_impl(Context, InteropProgram, nullptr) { - MIsInterop = true; -} - -program_impl::program_impl(ContextImplPtr Context, - pi_native_handle InteropProgram, - sycl::detail::pi::PiProgram Program) - : MProgram(Program), MContext(Context), MLinkable(true) { - const PluginPtr &Plugin = getPlugin(); - if (MProgram == nullptr) { - assert(InteropProgram && - "No InteropProgram/PiProgram defined with piextProgramFromNative"); - // Translate the raw program handle into PI program. - Plugin->call( - InteropProgram, MContext->getHandleRef(), false, &MProgram); - } else - Plugin->call(Program); - - // TODO handle the case when cl_program build is in progress - pi_uint32 NumDevices; - Plugin->call( - MProgram, PI_PROGRAM_INFO_NUM_DEVICES, sizeof(pi_uint32), &NumDevices, - nullptr); - std::vector PiDevices(NumDevices); - Plugin->call(MProgram, PI_PROGRAM_INFO_DEVICES, - sizeof(sycl::detail::pi::PiDevice) * - NumDevices, - PiDevices.data(), nullptr); - - std::vector PlatformDevices = - MContext->getPlatformImpl()->get_devices(); - // Keep only the subset of the devices (associated with context) that - // were actually used to create the program. - // This is possible when clCreateProgramWithBinary is used. - auto NewEnd = std::remove_if( - PlatformDevices.begin(), PlatformDevices.end(), - [&PiDevices](const sycl::device &Dev) { - return PiDevices.end() == - std::find(PiDevices.begin(), PiDevices.end(), - detail::getSyclObjImpl(Dev)->getHandleRef()); - }); - PlatformDevices.erase(NewEnd, PlatformDevices.end()); - MDevices = PlatformDevices; - assert(!MDevices.empty() && "No device found for this program"); - sycl::detail::pi::PiDevice Device = PiDevices[0]; - // TODO check build for each device instead - cl_program_binary_type BinaryType = PI_PROGRAM_BINARY_TYPE_NONE; - Plugin->call( - MProgram, Device, PI_PROGRAM_BUILD_INFO_BINARY_TYPE, - sizeof(cl_program_binary_type), &BinaryType, nullptr); - if (BinaryType == PI_PROGRAM_BINARY_TYPE_NONE) { - throw invalid_object_error( - "The native program passed to the program constructor has to be either " - "compiled or linked", - PI_ERROR_INVALID_PROGRAM); - } - size_t Size = 0; - Plugin->call( - MProgram, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, 0, nullptr, &Size); - std::vector OptionsVector(Size); - Plugin->call( - MProgram, Device, PI_PROGRAM_BUILD_INFO_OPTIONS, Size, - OptionsVector.data(), nullptr); - std::string Options(OptionsVector.begin(), OptionsVector.end()); - switch (BinaryType) { - case PI_PROGRAM_BINARY_TYPE_COMPILED_OBJECT: - MState = program_state::compiled; - MCompileOptions = Options; - MBuildOptions = Options; - return; - case PI_PROGRAM_BINARY_TYPE_LIBRARY: - case PI_PROGRAM_BINARY_TYPE_EXECUTABLE: - MState = program_state::linked; - MLinkOptions = ""; - MBuildOptions = Options; - return; - } - assert(false && "BinaryType is invalid."); -} - -program_impl::program_impl(ContextImplPtr Context, - sycl::detail::pi::PiKernel Kernel) - : program_impl(Context, reinterpret_cast(nullptr), - ProgramManager::getInstance().getPiProgramFromPiKernel( - Kernel, Context)) { - MIsInterop = true; -} - -program_impl::~program_impl() { - try { - // TODO catch an exception and put it to list of asynchronous exceptions - if (MProgram != nullptr) { - const PluginPtr &Plugin = getPlugin(); - Plugin->call(MProgram); - } - } catch (std::exception &e) { - __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~program_impl", e); - } -} - -cl_program program_impl::get() const { - throw_if_state_is(program_state::none); - getPlugin()->call(MProgram); - return pi::cast(MProgram); -} - -void program_impl::compile_with_kernel_name(std::string KernelName, - std::string CompileOptions) { - std::lock_guard Lock(MMutex); - throw_if_state_is_not(program_state::none); - create_pi_program_with_kernel_name( - KernelName, - /*JITCompilationIsRequired=*/(!CompileOptions.empty())); - compile(CompileOptions); - MState = program_state::compiled; -} - -void program_impl::link(std::string LinkOptions) { - std::lock_guard Lock(MMutex); - throw_if_state_is_not(program_state::compiled); - check_device_feature_support(MDevices); - std::vector Devices(get_pi_devices()); - const PluginPtr &Plugin = getPlugin(); - const char *LinkOpts = SYCLConfig::get(); - if (!LinkOpts) { - LinkOpts = LinkOptions.c_str(); - } - - // Plugin resets MProgram with a new pi_program as a result of the call to - // "piProgramLink". Thus, we need to release MProgram before the call to - // piProgramLink. - if (MProgram != nullptr) - Plugin->call(MProgram); - - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MContext->getHandleRef(), Devices.size(), Devices.data(), LinkOpts, - /*num_input_programs*/ 1, &MProgram, nullptr, nullptr, &MProgram); - Plugin->checkPiResult(Err); - MLinkOptions = LinkOptions; - MBuildOptions = LinkOptions; - MState = program_state::linked; -} - -bool program_impl::has_kernel(std::string KernelName, - bool /*IsCreatedFromSource*/) const { - throw_if_state_is(program_state::none); - - std::vector Devices(get_pi_devices()); - pi_uint64 function_ptr; - const PluginPtr &Plugin = getPlugin(); - - sycl::detail::pi::PiResult Err = PI_SUCCESS; - for (sycl::detail::pi::PiDevice Device : Devices) { - Err = Plugin->call_nocheck( - Device, MProgram, KernelName.c_str(), &function_ptr); - if (Err != PI_SUCCESS && - Err != PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE && - Err != PI_ERROR_INVALID_KERNEL_NAME) - throw runtime_error( - "Error from piextGetDeviceFunctionPointer when called by program", - Err); - if (Err == PI_SUCCESS || Err == PI_ERROR_FUNCTION_ADDRESS_IS_NOT_AVAILABLE) - return true; - } - - return false; -} - -kernel program_impl::get_kernel(std::string KernelName, - std::shared_ptr PtrToSelf, - bool IsCreatedFromSource) const { - throw_if_state_is(program_state::none); - auto [Kernel, ArgMask] = get_pi_kernel_arg_mask_pair(KernelName); - return createSyclObjFromImpl(std::make_shared( - Kernel, MContext, PtrToSelf, IsCreatedFromSource, nullptr, ArgMask)); -} - -std::vector> program_impl::get_binaries() const { - throw_if_state_is(program_state::none); - - std::vector> Result; - const PluginPtr &Plugin = getPlugin(); - std::vector BinarySizes(MDevices.size()); - Plugin->call( - MProgram, PI_PROGRAM_INFO_BINARY_SIZES, - sizeof(size_t) * BinarySizes.size(), BinarySizes.data(), nullptr); - - std::vector Pointers; - for (size_t I = 0; I < BinarySizes.size(); ++I) { - Result.emplace_back(BinarySizes[I]); - Pointers.push_back(Result[I].data()); - } - Plugin->call(MProgram, PI_PROGRAM_INFO_BINARIES, - sizeof(char *) * Pointers.size(), - Pointers.data(), nullptr); - return Result; -} - -void program_impl::compile(const std::string &Options) { - check_device_feature_support(MDevices); - std::vector Devices(get_pi_devices()); - const PluginPtr &Plugin = getPlugin(); - const char *CompileOpts = SYCLConfig::get(); - if (!CompileOpts) { - CompileOpts = Options.c_str(); - } - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MProgram, Devices.size(), Devices.data(), CompileOpts, 0, nullptr, - nullptr, nullptr, nullptr); - - if (Err != PI_SUCCESS) { - throw compile_program_error( - "Program compilation error:\n" + - ProgramManager::getProgramBuildLog(MProgram, MContext), - Err); - } - MCompileOptions = Options; - MBuildOptions = Options; -} - -void program_impl::build(const std::string &Options) { - check_device_feature_support(MDevices); - std::vector Devices(get_pi_devices()); - const PluginPtr &Plugin = getPlugin(); - ProgramManager::getInstance().flushSpecConstants(*this); - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MProgram, Devices.size(), Devices.data(), Options.c_str(), nullptr, - nullptr); - - if (Err != PI_SUCCESS) { - throw compile_program_error( - "Program build error:\n" + - ProgramManager::getProgramBuildLog(MProgram, MContext), - Err); - } - MBuildOptions = Options; -} - -std::vector program_impl::get_pi_devices() const { - std::vector PiDevices; - for (const auto &Device : MDevices) { - PiDevices.push_back(getSyclObjImpl(Device)->getHandleRef()); - } - return PiDevices; -} - -std::pair -program_impl::get_pi_kernel_arg_mask_pair(const std::string &KernelName) const { - std::pair Result; - - const PluginPtr &Plugin = getPlugin(); - sycl::detail::pi::PiResult Err = - Plugin->call_nocheck( - MProgram, KernelName.c_str(), &Result.first); - if (Err == PI_ERROR_INVALID_KERNEL_NAME) { - throw invalid_object_error( - "This instance of program does not contain the kernel requested", Err); - } - Plugin->checkPiResult(Err); - - // Some PI Plugins (like OpenCL) require this call to enable USM - // For others, PI will turn this into a NOP. - if (getContextImplPtr()->getPlatformImpl()->supports_usm()) - Plugin->call( - Result.first, PI_USM_INDIRECT_ACCESS, sizeof(pi_bool), &PI_TRUE); - - return Result; -} - -std::vector -program_impl::sort_devices_by_cl_device_id(std::vector Devices) { - std::sort(Devices.begin(), Devices.end(), - [](const device &id1, const device &id2) { - return (detail::getSyclObjImpl(id1)->getHandleRef() < - detail::getSyclObjImpl(id2)->getHandleRef()); - }); - return Devices; -} - -void program_impl::throw_if_state_is(program_state State) const { - if (MState == State) { - throw invalid_object_error("Invalid program state", - PI_ERROR_INVALID_PROGRAM); - } -} - -void program_impl::throw_if_state_is_not(program_state State) const { - if (MState != State) { - throw invalid_object_error("Invalid program state", - PI_ERROR_INVALID_PROGRAM); - } -} - -void program_impl::create_pi_program_with_kernel_name( - const std::string &KernelName, bool JITCompilationIsRequired) { - assert(!MProgram && "This program already has an encapsulated PI program"); - ProgramManager &PM = ProgramManager::getInstance(); - const device FirstDevice = get_devices()[0]; - RTDeviceBinaryImage &Img = PM.getDeviceImage( - KernelName, get_context(), FirstDevice, JITCompilationIsRequired); - MProgram = PM.createPIProgram(Img, get_context(), {FirstDevice}); -} - -void program_impl::flush_spec_constants( - const RTDeviceBinaryImage &Img, - sycl::detail::pi::PiProgram NativePrg) const { - // iterate via all specialization constants the program's image depends on, - // and set each to current runtime value (if any) - const RTDeviceBinaryImage::PropertyRange &SCRange = Img.getSpecConstants(); - ContextImplPtr Ctx = getSyclObjImpl(get_context()); - using SCItTy = RTDeviceBinaryImage::PropertyRange::ConstIterator; - - auto LockGuard = Ctx->getKernelProgramCache().acquireCachedPrograms(); - NativePrg = NativePrg ? NativePrg : getHandleRef(); - - for (SCItTy SCIt : SCRange) { - auto SCEntry = SpecConstRegistry.find((*SCIt)->Name); - if (SCEntry == SpecConstRegistry.end()) - // spec constant has not been set in user code - SPIR-V will use default - continue; - const spec_constant_impl &SC = SCEntry->second; - assert(SC.isSet() && "uninitialized spec constant"); - ByteArray Descriptors = DeviceBinaryProperty(*SCIt).asByteArray(); - - // First 8 bytes are consumed by the size of the property. - Descriptors.dropBytes(8); - - // Expected layout is vector of 3-component tuples (flattened into a - // vector of scalars), where each tuple consists of: ID of a scalar spec - // constant, (which might be a member of the composite); offset, which - // is used to calculate location of scalar member within the composite - // or zero for scalar spec constants; size of a spec constant. - while (!Descriptors.empty()) { - auto [Id, Offset, Size] = - Descriptors.consume(); - - Ctx->getPlugin()->call( - NativePrg, Id, Size, SC.getValuePtr() + Offset); - } - } -} - -pi_native_handle program_impl::getNative() const { - const auto &Plugin = getPlugin(); - if (getContextImplPtr()->getBackend() == backend::opencl) - Plugin->call(MProgram); - pi_native_handle Handle; - Plugin->call(MProgram, &Handle); - return Handle; -} - -} // namespace detail -} // namespace _V1 -} // namespace sycl diff --git a/sycl/source/detail/program_impl.hpp b/sycl/source/detail/program_impl.hpp deleted file mode 100644 index 67c02e95734ab..0000000000000 --- a/sycl/source/detail/program_impl.hpp +++ /dev/null @@ -1,396 +0,0 @@ -//==----- program_impl.hpp --- SYCL program implementation -----------------==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// -#pragma once - -#include -#include -#include -#include -#include -#include -#include - -#include -#include -#include -#include -#include - -namespace sycl { -inline namespace _V1 { - -// Forward declarations -class kernel; - -namespace detail { - -using ContextImplPtr = std::shared_ptr; - -enum class program_state { none = 0, compiled = 1, linked = 2 }; - -class program_impl { -public: - program_impl() = delete; - - /// Constructs an instance of program. - /// - /// The program will be created in the program_state::none state and - /// associated with the provided context and the devices that are associated - /// with the context. - /// - /// \param Context is a pointer to SYCL context impl. - /// \param PropList is an instance of property_list. - explicit program_impl(ContextImplPtr Context, const property_list &PropList); - - /// Constructs an instance of SYCL program for the provided DeviceList. - /// - /// The program will be created in the program_state::none state and - /// associated with the provided context and the devices in the provided - /// DeviceList. - /// - /// \param Context is a pointer to SYCL context impl. - /// \param DeviceList is a list of SYCL devices. - /// \param PropList is an instance of property_list. - program_impl(ContextImplPtr Context, std::vector DeviceList, - const property_list &PropList); - - /// Constructs an instance of SYCL program by linking together each SYCL - /// program instance in ProgramList. - /// - /// Each program in ProgramList must be in the program_state::compiled - /// state and must be associated with the same SYCL context. Otherwise an - /// invalid_object_error SYCL exception will be thrown. A - /// feature_not_supported exception will be thrown if any device that the - /// program is to be linked for returns false for the device information - /// query info::device::is_linker_available. Kernels caching for linked - /// programs won't be allowed due to only compiled state of each and every - /// program in the list and thus unknown state of caching resolution. - /// - /// \param ProgramList is a list of program_impl instances. - /// \param LinkOptions is a string containing valid OpenCL link options. - /// \param PropList is an instance of property_list. - program_impl(std::vector> ProgramList, - std::string LinkOptions, const property_list &PropList); - - /// Constructs a program instance from an interop raw BE program handle. - /// TODO: BE generalization will change that to something better. - /// - /// The state of the constructed program can be either - /// program_state::compiled or program_state::linked, depending on the state - /// of the InteropProgram. Otherwise an invalid_object_error SYCL exception is - /// thrown. - /// - /// The instance of the program will be retained on construction. - /// - /// \param Context is a pointer to SYCL context impl. - /// \param InteropProgram is an instance of plugin interface interoperability - /// program. - program_impl(ContextImplPtr Context, pi_native_handle InteropProgram); - - /// Constructs a program instance from plugin interface interoperability - /// kernel. - /// - /// \param Context is a pointer to SYCL context impl. - /// \param Kernel is a raw PI kernel handle. - program_impl(ContextImplPtr Context, sycl::detail::pi::PiKernel Kernel); - - ~program_impl(); - - /// Checks if this program_impl has a property of type propertyT. - /// - /// \return true if this program_impl has a property of type propertyT. - template bool has_property() const { - return MPropList.has_property(); - } - - /// Gets the specified property of this program_impl. - /// - /// Throws invalid_object_error if this program_impl does not have a property - /// of type propertyT. - /// - /// \return a copy of the property of type propertyT. - template propertyT get_property() const { - return MPropList.get_property(); - } - - /// Returns a valid cl_program instance. - /// - /// The instance of cl_program will be retained before returning. - /// If the program is created for a SYCL host device, an - /// invalid_object_error exception is thrown. - /// - /// \return a valid OpenCL cl_program instance. - cl_program get() const; - - /// \return a reference to a raw PI program handle. PI program is not - /// retained before return. - sycl::detail::pi::PiProgram &getHandleRef() { return MProgram; } - /// \return a constant reference to a raw PI program handle. PI program is - /// not retained before return. - const sycl::detail::pi::PiProgram &getHandleRef() const { return MProgram; } - - /// Compiles the SYCL kernel function into the encapsulated raw program. - /// - /// The kernel function is defined by its name. This member function - /// sets the state of this SYCL program to program_state::compiled. - /// If this program was not in the program_state::none state, - /// an invalid_object_error exception is thrown. If the compilation fails, - /// a compile_program_error SYCL exception is thrown. If any device that the - /// program is being compiled for returns false for the device information - /// query info::device::is_compiler_available, a feature_not_supported - /// exception is thrown. - /// - /// \param KernelName is a string containing SYCL kernel name. - /// \param CompileOptions is a string of valid OpenCL compile options. - /// \param Module is an OS handle to user code module. - void compile_with_kernel_name(std::string KernelName, - std::string CompileOptions); - - /// Builds the SYCL kernel function into encapsulated raw program. - /// - /// The SYCL kernel function is defined by the kernel name. - /// This member function sets the state of this SYCL program to - /// program_state::linked. If the program was not in the program_state::none - /// state, an invalid_object_error SYCL exception is thrown. If the - /// compilation fails, a compile_program_error SYCL exception is thrown. If - /// any device that the program is being built for returns false for the - /// device information queries info::device::is_compiler_available or - /// info::device::is_linker_available, a feature_not_supported SYCL - /// exception is thrown. - /// - /// \param KernelName is a string containing SYCL kernel name. - /// \param BuildOptions is a string containing OpenCL compile options. - /// \param M is an OS handle to user code module. - void build_with_kernel_name(std::string KernelName, std::string BuildOptions); - - /// Links encapsulated raw program. - /// - /// This member function sets the state of this SYCL program to - /// program_state::linked. If the program was not in the - /// program_state::compiled state, an invalid_object_error SYCL exception is - /// thrown. If linking fails, a compile_program_error is thrown. If any - /// device that the program is to be linked for returns false for the device - /// information query info::device::is_linker_available, a - /// feature_not_supported exception is thrown. - /// - /// \param LinkOptions is a string containing OpenCL link options. - void link(std::string LinkOptions = ""); - - /// Checks if kernel is available for this program. - /// - /// The SYCL kernel is defined by kernel name. If the program state is - /// program_state::none an invalid_object_error SYCL exception is thrown. - /// - /// \return true if the SYCL kernel is available. - bool has_kernel(std::string KernelName, bool IsCreatedFromSource) const; - - /// Returns a SYCL kernel for the SYCL kernel function defined by kernel - /// name. - /// - /// If program is in the program_state::none state or if the SYCL kernel - /// function is not available, an invalid_object_error exception is thrown. - /// - /// \return a valid instance of SYCL kernel. - kernel get_kernel(std::string KernelName, - std::shared_ptr PtrToSelf, - bool IsCreatedFromSource) const; - - /// Returns built program binaries. - /// - /// If this program is not in the program_state::compiled or - /// program_state::linked states, an invalid_object_error SYCL exception - /// is thrown. - /// - /// \return a vector of vectors representing the compiled binaries for each - /// associated SYCL device. - std::vector> get_binaries() const; - - /// \return the SYCL context that this program was constructed with. - context get_context() const { - return createSyclObjFromImpl(MContext); - } - - /// \return the Plugin associated with the context of this program. - const PluginPtr &getPlugin() const { return MContext->getPlugin(); } - - ContextImplPtr getContextImplPtr() const { return MContext; } - - /// \return a vector of devices that are associated with this program. - std::vector get_devices() const { return MDevices; } - - /// Returns compile options that were provided when the encapsulated program - /// was explicitly compiled. - /// - /// If the program was built instead of explicitly compiled, if the program - /// has not yet been compiled, or if the program has been compiled for only - /// the host device, then an empty string is return, unless the underlying - /// cl_program was explicitly compiled, in which case the compile options - /// used in the explicit compile are returned. - /// - /// \return a string of valid OpenCL compile options. - std::string get_compile_options() const { return MCompileOptions; } - - /// Returns compile options that were provided to the most recent invocation - /// of link member function. - /// - /// If the program has not been explicitly linked using the aforementioned - /// function, constructed with an explicitly linking constructor, or if the - /// program has been linked for only the host device, then an empty string - /// is returned. If the program was constructed from cl_program, then an - /// empty string is returned unless the cl_program was explicitly linked, - /// in which case the link options used in that explicit link are returned. - /// If the program object was constructed using a constructor form that - /// links a vector of programs, then the link options passed to this - /// constructor are returned. - /// - /// \return a string of valid OpenCL compile options. - std::string get_link_options() const { return MLinkOptions; } - - /// Returns the compile, link, or build options, from whichever of those - /// operations was performed most recently on the encapsulated cl_program. - /// - /// If no compile, link, or build operations have been performed on this - /// program, or if the program includes the host device in its device list, - /// then an empty string is returned. - /// - /// \return a string of valid OpenCL build options. - std::string get_build_options() const { return MBuildOptions; } - - /// \return the current state of this SYCL program. - program_state get_state() const { return MState; } - - /// Takes current values of specialization constants and "injects" them into - /// the underlying native program program via specialization constant - /// managemment PI APIs. The native program passed as non-null argument - /// overrides the MProgram native program field. - /// \param Img device binary image corresponding to this program, used to - /// resolve spec constant name to SPIR-V integer ID - /// \param NativePrg if not null, used as the flush target, otherwise MProgram - /// is used - void - flush_spec_constants(const RTDeviceBinaryImage &Img, - sycl::detail::pi::PiProgram NativePrg = nullptr) const; - - void stableSerializeSpecConstRegistry(SerializedObj &Dst) const { - detail::stableSerializeSpecConstRegistry(SpecConstRegistry, Dst); - } - - /// Tells whether a specialization constant has been set for this program. - bool hasSetSpecConstants() const { return !SpecConstRegistry.empty(); } - - /// \return true if caching is allowed for this program. - bool is_cacheable() const { return MProgramAndKernelCachingAllowed; } - - /// Returns the native plugin handle. - pi_native_handle getNative() const; - - bool isInterop() const { return MIsInterop; } - -private: - // Deligating Constructor used in Implementation. - program_impl(ContextImplPtr Context, pi_native_handle InteropProgram, - sycl::detail::pi::PiProgram Program); - /// Checks feature support for specific devices. - /// - /// If there's at least one device that does not support this feature, - /// a feature_not_supported exception is thrown. - /// - /// \param Devices is a vector of SYCL devices. - template - void check_device_feature_support(const std::vector &Devices) { - for (const auto &Device : Devices) { - if (!Device.get_info()) { - throw sycl::exception( - sycl::errc::feature_not_supported, - "Online compilation is not supported by this device"); - } - } - } - - /// Creates a plugin interface kernel using its name. - /// - /// \param Module is an OS handle to user code module. - /// \param KernelName is a name of kernel to be created. - /// \param JITCompilationIsRequired If JITCompilationIsRequired is true - /// add a check that kernel is compiled, otherwise don't add the check. - void - create_pi_program_with_kernel_name(const std::string &KernelName, - bool JITCompilationIsRequired = false); - - /// Compiles underlying plugin interface program. - /// - /// \param Options is a string containing OpenCL compile options. - void compile(const std::string &Options); - - /// Builds underlying plugin interface program. - /// - /// \param Options is a string containing OpenCL build options. - void build(const std::string &Options); - - /// \return a vector of devices managed by the plugin. - std::vector get_pi_devices() const; - - /// \param Options is a string containing OpenCL C build options. - /// \return true if caching is allowed for this program and build options. - static bool is_cacheable_with_options(const std::string &Options) { - return Options.empty(); - } - - /// \param KernelName is a string containing OpenCL kernel name. - /// \return true if underlying OpenCL program has kernel with specific name. - bool has_cl_kernel(const std::string &KernelName) const; - - /// \param KernelName is a string containing PI kernel name. - /// \return an instance of PI kernel with specific name. If kernel is - /// unavailable, an invalid_object_error exception is thrown. - std::pair - get_pi_kernel_arg_mask_pair(const std::string &KernelName) const; - - /// \return a vector of sorted in ascending order SYCL devices. - std::vector sort_devices_by_cl_device_id(std::vector Devices); - - /// Throws an invalid_object_exception if state of this program is in the - /// specified state. - /// - /// \param State is a program state to match against. - void throw_if_state_is(program_state State) const; - - /// Throws an invalid_object_exception if state of this program is not in - /// the specified state. - /// - /// \param State is a program state to match against. - void throw_if_state_is_not(program_state State) const; - - sycl::detail::pi::PiProgram MProgram = nullptr; - program_state MState = program_state::none; - std::mutex MMutex; - ContextImplPtr MContext; - bool MLinkable = false; - std::vector MDevices; - property_list MPropList; - std::string MCompileOptions; - std::string MLinkOptions; - std::string MBuildOptions; - - // Keeps specialization constant map for this program. Spec constant name - // resolution to actual SPIR-V integer ID happens at build time, where the - // device binary image is available. Access is guarded by this context's - // program cache lock. - SpecConstRegistryT SpecConstRegistry; - - /// Only allow kernel caching for programs constructed with context only (or - /// device list and context) and built with build_with_kernel_type with - /// default build options - bool MProgramAndKernelCachingAllowed = false; - - bool MIsInterop = false; -}; - -} // namespace detail -} // namespace _V1 -} // namespace sycl diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 82246af25173d..dfec4a811382f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -14,7 +14,6 @@ #include #include #include -#include #include #include #include @@ -1452,45 +1451,6 @@ void ProgramManager::dumpImage(const RTDeviceBinaryImage &Img, F.close(); } -void ProgramManager::flushSpecConstants(const program_impl &Prg, - sycl::detail::pi::PiProgram NativePrg, - const RTDeviceBinaryImage *Img) { - if (DbgProgMgr > 2) { - std::cerr << ">>> ProgramManager::flushSpecConstants(" << Prg.get() - << ",...)\n"; - } - if (!Prg.hasSetSpecConstants()) - return; // nothing to do - pi::PiProgram PrgHandle = Prg.getHandleRef(); - // program_impl can't correspond to two different native programs - assert(!NativePrg || !PrgHandle || (NativePrg == PrgHandle)); - NativePrg = NativePrg ? NativePrg : PrgHandle; - - if (!Img) { - // caller hasn't provided the image object - find it - { // make sure NativePrograms map access is synchronized - std::lock_guard Lock(MNativeProgramsMutex); - auto It = NativePrograms.find(NativePrg); - if (It == NativePrograms.end()) - throw sycl::exception( - sycl::errc::invalid, - "spec constant is set in a program w/o a binary image"); - Img = It->second; - } - if (!Img->supportsSpecConstants()) { - if (DbgProgMgr > 0) - std::cerr << ">>> ProgramManager::flushSpecConstants: binary image " - << &Img->getRawData() << " doesn't support spec constants\n"; - // This device binary image does not support runtime setting of - // specialization constants; compiler must have generated default values. - // NOTE: Can't throw here, as it would always take place with AOT - //-compiled code. New Khronos 2020 spec should fix this inconsistency. - return; - } - } - Prg.flush_spec_constants(*Img, NativePrg); -} - uint32_t ProgramManager::getDeviceLibReqMask(const RTDeviceBinaryImage &Img) { const RTDeviceBinaryImage::PropertyRange &DLMRange = Img.getDeviceLibReqMask(); diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 573e4ddfed284..20a2e92406f21 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -64,7 +64,6 @@ class context_impl; using ContextImplPtr = std::shared_ptr; class device_impl; using DeviceImplPtr = std::shared_ptr; -class program_impl; class queue_impl; class event_impl; // DeviceLibExt is shared between sycl runtime and sycl-post-link tool. @@ -165,20 +164,6 @@ class ProgramManager { getProgramBuildLog(const sycl::detail::pi::PiProgram &Program, const ContextImplPtr Context); - /// Resolves given program to a device binary image and requests the program - /// to flush constants the image depends on. - /// \param Prg the program object to get spec constant settings from. - /// Passing program_impl by raw reference is OK, since it is not - /// captured anywhere once the function returns. - /// \param NativePrg the native program, target for spec constant setting; if - /// not null then overrides the native program in Prg - /// \param Img A source of the information about which constants need - /// setting and symboling->integer spec constant ID mapping. If not - /// null, overrides native program->binary image binding maintained by - /// the program manager. - void flushSpecConstants(const program_impl &Prg, - pi::PiProgram NativePrg = nullptr, - const RTDeviceBinaryImage *Img = nullptr); uint32_t getDeviceLibReqMask(const RTDeviceBinaryImage &Img); /// Returns the mask for eliminated kernel arguments for the requested kernel diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index b9ff259906f95..b9b2fd6c5b280 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -15,7 +15,6 @@ #include #include #include -#include #include #include #include diff --git a/sycl/unittests/kernel-and-program/Cache.cpp b/sycl/unittests/kernel-and-program/Cache.cpp index f29e75299bbd2..7205c1eaea867 100644 --- a/sycl/unittests/kernel-and-program/Cache.cpp +++ b/sycl/unittests/kernel-and-program/Cache.cpp @@ -13,7 +13,6 @@ #include "detail/context_impl.hpp" #include "detail/kernel_program_cache.hpp" -#include "detail/program_impl.hpp" #include "sycl/detail/pi.h" #include #include