diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index bd7fac34d61b..7fd27f75c062 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -1101,8 +1101,6 @@ struct is_property_key_of make_kernel_bundle_from_source( diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index f2f2673562f8..679bfa407ee2 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -891,8 +891,12 @@ bool device_impl::isGetDeviceAndHostTimerSupported() { bool device_impl::extOneapiCanCompile( ext::oneapi::experimental::source_language Language) { try { + // Get the shared_ptr to this object from the platform that owns it. + std::shared_ptr Self = MPlatform->getOrMakeDeviceImpl(MDevice); return sycl::ext::oneapi::experimental::detail:: - is_source_kernel_bundle_supported(getBackend(), Language); + is_source_kernel_bundle_supported(Language, + std::vector{Self}); + } catch (sycl::exception &) { return false; } diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 12ac9a47b7c0..94744b4813e1 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -30,6 +31,18 @@ namespace sycl { inline namespace _V1 { + +namespace ext::oneapi::experimental::detail { +using DeviceImplPtr = std::shared_ptr; +bool is_source_kernel_bundle_supported( + sycl::ext::oneapi::experimental::source_language Language, + const context &Ctx); + +bool is_source_kernel_bundle_supported( + sycl::ext::oneapi::experimental::source_language Language, + const std::vector &Devices); +} // namespace ext::oneapi::experimental::detail + namespace detail { static bool checkAllDevicesAreInContext(const std::vector &Devices, diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index 63907ff913dc..78e1f1399a2c 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -25,6 +25,10 @@ inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { +// forward declaration +std::string InvokeOclocQuery(const std::vector &IPVersionVec, + const char *identifier); + // ensures the OclocLibrary has the right version, etc. void checkOclocLibrary(void *OclocLibrary) { void *OclocVersionHandle = @@ -64,35 +68,40 @@ static std::unique_ptr> std::ignore = sycl::detail::ur::unloadOsLibrary(StoredPtr); }); -// load the ocloc shared library, check it. -void loadOclocLibrary() { +void loadOclocLibrary(const std::vector &IPVersionVec) { #ifdef __SYCL_RT_OS_WINDOWS - static const std::string OclocLibraryName = "ocloc64.dll"; + static const std::string OclocPath = "ocloc64.dll"; #else - static const std::string OclocLibraryName = "libocloc.so"; + static const std::string OclocPath = "libocloc.so"; #endif - void *tempPtr = OclocLibrary.get(); - if (tempPtr == nullptr) { - tempPtr = sycl::detail::ur::loadOsLibrary(OclocLibraryName); + + // set OclocLibrary value by side effect. + try { + // Load then perform checks. Each check throws. + void *tempPtr = sycl::detail::ur::loadOsLibrary(OclocPath); + OclocLibrary.reset(tempPtr); if (tempPtr == nullptr) throw sycl::exception(make_error_code(errc::build), - "Unable to load ocloc library " + OclocLibraryName); + "Unable to load ocloc from " + OclocPath); checkOclocLibrary(tempPtr); - OclocLibrary.reset(tempPtr); + InvokeOclocQuery(IPVersionVec, "CL_DEVICE_OPENCL_C_ALL_VERSIONS"); + } catch (const sycl::exception &) { + OclocLibrary.reset(nullptr); + std::rethrow_exception(std::current_exception()); } } -bool OpenCLC_Compilation_Available() { +bool OpenCLC_Compilation_Available(const std::vector &IPVersionVec) { // Already loaded? if (OclocLibrary != nullptr) return true; try { // loads and checks version - loadOclocLibrary(); + loadOclocLibrary(IPVersionVec); return true; } catch (...) { return false; @@ -102,11 +111,12 @@ bool OpenCLC_Compilation_Available() { using voidPtr = void *; void SetupLibrary(voidPtr &oclocInvokeHandle, voidPtr &oclocFreeOutputHandle, - std::error_code the_errc) { - if (!oclocInvokeHandle) { - if (OclocLibrary == nullptr) - loadOclocLibrary(); + std::error_code the_errc, + const std::vector &IPVersionVec) { + if (OclocLibrary == nullptr) + loadOclocLibrary(IPVersionVec); + if (!oclocInvokeHandle) { oclocInvokeHandle = sycl::detail::ur::getOsLibraryFuncAddress( OclocLibrary.get(), "oclocInvoke"); if (!oclocInvokeHandle) @@ -145,7 +155,8 @@ std::string InvokeOclocQuery(const std::vector &IPVersionVec, static void *oclocFreeOutputHandle = nullptr; std::error_code the_errc = make_error_code(errc::runtime); - SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, the_errc); + SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, the_errc, + IPVersionVec); uint32_t NumOutputs = 0; uint8_t **Outputs = nullptr; @@ -205,7 +216,8 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, static void *oclocFreeOutputHandle = nullptr; std::error_code build_errc = make_error_code(errc::build); - SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, build_errc); + SetupLibrary(oclocInvokeHandle, oclocFreeOutputHandle, build_errc, + IPVersionVec); // assemble ocloc args std::string CombinedUserArgs = diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index d618c86e07d9..85f4b9b2a488 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -24,8 +24,8 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &IPVersionVec, const std::vector &UserArgs, std::string *LogPtr); - -bool OpenCLC_Compilation_Available(); +// IPVersionVec gets flattened and passed to ocloc as the -dev flag. +bool OpenCLC_Compilation_Available(const std::vector &IPVersionVec); bool OpenCLC_Feature_Available(const std::string &Feature, uint32_t IPVersion); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index ee3dbc78319e..eee4fc23cfb2 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -387,24 +387,62 @@ namespace detail { ///////////////////////// // syclex::detail::is_source_kernel_bundle_supported ///////////////////////// -bool is_source_kernel_bundle_supported(backend BE, source_language Language) { + +bool is_source_kernel_bundle_supported( + sycl::ext::oneapi::experimental::source_language Language, + const std::vector &DeviceImplVec) { + backend BE = DeviceImplVec[0]->getBackend(); // Support is limited to the opencl and level_zero backends. bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) || (BE == sycl::backend::opencl); - if (BE_Acceptable) { - if (Language == source_language::opencl) { - return detail::OpenCLC_Compilation_Available(); - } else if (Language == source_language::spirv) { - return true; - } else if (Language == source_language::sycl) { - return detail::SYCL_JIT_Compilation_Available(); - } + if (!BE_Acceptable) + return false; + + if (Language == source_language::spirv) { + return true; + } else if (Language == source_language::sycl) { + return detail::SYCL_JIT_Compilation_Available(); + } else if (Language == source_language::opencl) { + if (DeviceImplVec.empty()) + return false; + + const AdapterPtr &Adapter = DeviceImplVec[0]->getAdapter(); + std::vector IPVersionVec; + IPVersionVec.reserve(DeviceImplVec.size()); + + std::transform(DeviceImplVec.begin(), DeviceImplVec.end(), + std::back_inserter(IPVersionVec), + [&](const DeviceImplPtr &Impl) { + uint32_t ipVersion = 0; + ur_device_handle_t DeviceHandle = Impl->getHandleRef(); + Adapter->call( + DeviceHandle, UR_DEVICE_INFO_IP_VERSION, + sizeof(uint32_t), &ipVersion, nullptr); + return ipVersion; + }); + + return detail::OpenCLC_Compilation_Available(IPVersionVec); } // otherwise return false; } +bool is_source_kernel_bundle_supported( + sycl::ext::oneapi::experimental::source_language Language, + const context &Ctx) { + const std::vector Devices = Ctx.get_devices(); + std::vector DeviceImplVec; + DeviceImplVec.reserve(Devices.size()); + std::transform(Devices.begin(), Devices.end(), + std::back_inserter(DeviceImplVec), + [](const sycl::device &dev) { + return sycl::detail::getSyclObjImpl(dev); + }); + + return is_source_kernel_bundle_supported(Language, DeviceImplVec); +} + ///////////////////////// // syclex::detail::create_kernel_bundle_from_source ///////////////////////// @@ -428,8 +466,7 @@ make_kernel_bundle_from_source(const context &SyclContext, for (auto &p : IncludePairViews) IncludePairs.push_back({p.first.data(), p.second.data()}); - backend BE = SyclContext.get_backend(); - if (!is_source_kernel_bundle_supported(BE, Language)) + if (!is_source_kernel_bundle_supported(Language, SyclContext)) throw sycl::exception(make_error_code(errc::invalid), "kernel_bundle creation from source not supported"); @@ -448,8 +485,7 @@ source_kb make_kernel_bundle_from_source(const context &SyclContext, const std::vector &Bytes, include_pairs_view_t IncludePairs) { (void)IncludePairs; - backend BE = SyclContext.get_backend(); - if (!is_source_kernel_bundle_supported(BE, Language)) + if (!is_source_kernel_bundle_supported(Language, SyclContext)) throw sycl::exception(make_error_code(errc::invalid), "kernel_bundle creation from source not supported"); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 967e13b09d73..d27ff51778f7 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3105,7 +3105,6 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_5 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageENS0_6detail11string_viewESt6vectorISt4pairISA_SA_ESaISD_EE _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaISA_EES9_ISt4pairINS0_6detail11string_viewESH_ESaISI_EE -_ZN4sycl3_V13ext6oneapi12experimental6detail33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental6memcpyENS0_5queueEPvPKvmRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental6memsetENS0_5queueEPvimRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 27a22a704432..dd3e8c2e80d6 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4226,7 +4226,6 @@ ?is_gpu@device@_V1@sycl@@QEBA_NXZ ?is_in_fusion_mode@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEBA_NXZ ?is_in_order@queue@_V1@sycl@@QEBA_NXZ -?is_source_kernel_bundle_supported@detail@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@56@W4source_language@23456@@Z ?is_specialization_constant_set@kernel_bundle_plain@detail@_V1@sycl@@IEBA_NPEBD@Z ?join_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@2@@5@W4bundle_state@23@@Z ?khr_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ