diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc similarity index 97% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index 076cd07df6add..09a6fe5645fcc 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -49,14 +49,12 @@ This extension also depends on the following other SYCL extensions: == Status -This is a proposed extension specification, intended to gather community -feedback. -Interfaces defined in this specification may not be implemented yet or may be -in a preliminary state. -The specification itself may also change in incompatible ways before it is -finalized. -*Shipping software products should not rely on APIs defined in this -specification.* +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.* == Overview diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc similarity index 96% rename from sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc rename to sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc index f6c0b3acf1a19..5322338e3aac1 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc @@ -49,14 +49,12 @@ This extension also depends on the following other SYCL extensions: == Status -This is a proposed extension specification, intended to gather community -feedback. -Interfaces defined in this specification may not be implemented yet or may be -in a preliminary state. -The specification itself may also change in incompatible ways before it is -finalized. -*Shipping software products should not rely on APIs defined in this -specification.* +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.* == Overview diff --git a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp index ebb51028f8d99..15536e9300b90 100644 --- a/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp +++ b/sycl/include/sycl/ext/intel/experimental/online_compiler.hpp @@ -75,7 +75,11 @@ enum class source_language { opencl_c = 0, cm = 1 }; /// Represents an online compiler for the language given as template /// parameter. -template class online_compiler { +template +class __SYCL2020_DEPRECATED( + "experimental online_compiler is being deprecated. See " + "'sycl_ext_oneapi_kernel_compiler.asciidoc' instead for new kernel " + "compiler extension to kernel_bundle implementation.") online_compiler { public: /// Constructs online compiler which can target any device and produces /// given compiled code format. Produces 64-bit device code. diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index 2fcc5f7004a5d..d3cca173c630f 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -241,6 +241,30 @@ struct ValueOrDefault< } }; +template struct all_props_are_keys_of; + +template +struct all_props_are_keys_of : std::true_type {}; + +template +struct all_props_are_keys_of< + SyclT, ext::oneapi::experimental::detail::empty_properties_t> + : std::true_type {}; + +template +struct all_props_are_keys_of< + SyclT, ext::oneapi::experimental::properties>> + : std::bool_constant< + ext::oneapi::experimental::is_property_key_of::value> { +}; + +template +struct all_props_are_keys_of< + SyclT, ext::oneapi::experimental::properties>> + : std::bool_constant< + ext::oneapi::experimental::is_property_key_of::value && + all_props_are_keys_of()> {}; + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 304356f71b30f..45b326399b7dc 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -218,8 +218,10 @@ enum PropKind : uint32_t { CacheControlReadHint = 48, CacheControlReadAssertion = 49, CacheControlWrite = 50, + BuildOptions = 51, + BuildLog = 52, // PropKindSize must always be the last value. - PropKindSize = 51, + PropKindSize = 53, }; // This trait must be specialized for all properties and must have a unique diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index fdd2b0c77bd1b..b2893854a79d4 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -20,6 +20,10 @@ #include // for bundle_state #include // for property_list +#include // PropertyT +#include // build_options +#include // and log + #include // for array #include // for size_t, memcpy #include // for function @@ -41,6 +45,7 @@ auto get_native(const kernel_bundle &Obj) namespace detail { class kernel_id_impl; +class kernel_impl; } template kernel_id get_kernel_id(); @@ -176,6 +181,10 @@ class __SYCL_EXPORT kernel_bundle_plain { bool native_specialization_constant() const noexcept; + bool ext_oneapi_has_kernel(const std::string &name); + + kernel ext_oneapi_get_kernel(const std::string &name); + protected: // \returns a kernel object which represents the kernel identified by // kernel_id passed @@ -215,7 +224,12 @@ class kernel_bundle : public detail::kernel_bundle_plain, kernel_bundle() = delete; /// \returns true if the kernel_bundles contains no device images - bool empty() const noexcept { return kernel_bundle_plain::empty(); } + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> + bool empty() const noexcept { + return kernel_bundle_plain::empty(); + } /// \returns the backend associated with the kernel bundle backend get_backend() const noexcept { @@ -234,6 +248,9 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \returns true if the kernel_bundle contains the kernel identified by /// kernel_id passed + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> bool has_kernel(const kernel_id &KernelID) const noexcept { return kernel_bundle_plain::has_kernel(KernelID); } @@ -241,36 +258,53 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \returns true if the kernel_bundle contains the kernel identified by /// kernel_id passed and if this kernel is compatible with the device /// specified + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> bool has_kernel(const kernel_id &KernelID, const device &Dev) const noexcept { return kernel_bundle_plain::has_kernel(KernelID, Dev); } /// \returns true only if the kernel bundle contains the kernel identified by /// KernelName. - template bool has_kernel() const noexcept { + template < + typename KernelName, bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> + bool has_kernel() const noexcept { return has_kernel(get_kernel_id()); } /// \returns true only if the kernel bundle contains the kernel identified by /// KernelName and if that kernel is compatible with the device Dev. - template + template < + typename KernelName, bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> bool has_kernel(const device &Dev) const noexcept { return has_kernel(get_kernel_id(), Dev); } /// \returns a vector of kernel_id's that contained in the kernel_bundle + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> std::vector get_kernel_ids() const { return kernel_bundle_plain::get_kernel_ids(); } /// \returns true if the kernel_bundle contains at least one device image /// which uses specialization constants + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> bool contains_specialization_constants() const noexcept { return kernel_bundle_plain::contains_specialization_constants(); } /// \returns true if all specialization constants which are used in the /// kernel_bundle are "native specialization constants in all device images + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> bool native_specialization_constant() const noexcept { return kernel_bundle_plain::native_specialization_constant(); } @@ -293,7 +327,10 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \returns true if any device image in the kernel_bundle uses specialization /// constant whose address is SpecName - template bool has_specialization_constant() const noexcept { + template < + auto &SpecName, bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> + bool has_specialization_constant() const noexcept { const char *SpecSymName = detail::get_spec_constant_symbolic_ID(); return has_specialization_constant_impl(SpecSymName); } @@ -312,7 +349,9 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \returns the value of the specialization constant whose address is /// SpecName for this kernel bundle. - template + template < + auto &SpecName, bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> typename std::remove_reference_t::value_type get_specialization_constant() const { using SCType = @@ -331,16 +370,42 @@ class kernel_bundle : public detail::kernel_bundle_plain, } /// \returns an iterator to the first device image kernel_bundle contains + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> device_image_iterator begin() const { return reinterpret_cast( kernel_bundle_plain::begin()); } /// \returns an iterator to the last device image kernel_bundle contains + template < + bundle_state _State = State, + typename = std::enable_if_t<_State != bundle_state::ext_oneapi_source>> device_image_iterator end() const { return reinterpret_cast(kernel_bundle_plain::end()); } + ///////////////////////// + // ext_oneapi_has_kernel + // only true if created from source and has this kernel + ///////////////////////// + template > + bool ext_oneapi_has_kernel(const std::string &name) { + return detail::kernel_bundle_plain::ext_oneapi_has_kernel(name); + } + + ///////////////////////// + // ext_oneapi_get_kernel + // kernel_bundle must be created from source, throws if not present + ///////////////////////// + template > + kernel ext_oneapi_get_kernel(const std::string &name) { + return detail::kernel_bundle_plain::ext_oneapi_get_kernel(name); + } + private: kernel_bundle(detail::KernelBundleImplPtr Impl) : kernel_bundle_plain(std::move(Impl)) {} @@ -743,6 +808,138 @@ build(const kernel_bundle &InputBundle, return build(InputBundle, InputBundle.get_devices(), PropList); } +namespace ext::oneapi::experimental { + +///////////////////////// +// PropertyT syclex::build_options +///////////////////////// +struct build_options { + std::vector opts; + build_options(const std::string &optsArg) : opts{optsArg} {} + build_options(const std::vector &optsArg) : opts(optsArg) {} +}; +using build_options_key = build_options; + +template <> struct is_property_key : std::true_type {}; + +template <> +struct is_property_key_of> + : std::true_type {}; + +namespace detail { + +template <> +struct PropertyToKind { + static constexpr PropKind Kind = PropKind::BuildOptions; +}; + +template <> +struct IsRuntimeProperty + : std::true_type {}; + +template <> +struct IsCompileTimeProperty + : std::false_type {}; + +} // namespace detail + +///////////////////////// +// PropertyT syclex::save_log +///////////////////////// +struct save_log { + std::string *log; + save_log(std::string *logArg) : log(logArg) {} +}; +using save_log_key = save_log; + +template <> struct is_property_key : std::true_type {}; + +template <> +struct is_property_key_of> + : std::true_type {}; + +namespace detail { + +template <> +struct PropertyToKind { + static constexpr PropKind Kind = PropKind::BuildLog; +}; + +template <> +struct IsRuntimeProperty + : std::true_type {}; + +template <> +struct IsCompileTimeProperty + : std::false_type {}; + +} // namespace detail + +///////////////////////// +// syclex::is_source_kernel_bundle_supported +///////////////////////// +__SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, + source_language Language); + +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +__SYCL_EXPORT kernel_bundle +create_kernel_bundle_from_source( + const context &SyclContext, + sycl::ext::oneapi::experimental::source_language Language, + const std::string &Source); + +///////////////////////// +// syclex::build(source_kb) => exe_kb +///////////////////////// +namespace detail { +// forward decl +__SYCL_EXPORT kernel_bundle +build_from_source(kernel_bundle &SourceKB, + const std::vector &Devices, + const std::vector &BuildOptions, + std::string *LogPtr); + +} // namespace detail + +template && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> + +kernel_bundle +build(kernel_bundle &SourceKB, + const std::vector &Devices, PropertyListT props = {}) { + std::vector BuildOptionsVec; + std::string *LogPtr = nullptr; + if constexpr (props.template has_property()) { + BuildOptionsVec = props.template get_property().opts; + } + if constexpr (props.template has_property()) { + LogPtr = props.template get_property().log; + } + return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); +} + +template && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> +kernel_bundle +build(kernel_bundle &SourceKB, + PropertyListT props = {}) { + return build(SourceKB, SourceKB.get_devices(), props); +} + +} // namespace ext::oneapi::experimental + } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/kernel_bundle_enums.hpp b/sycl/include/sycl/kernel_bundle_enums.hpp index 45a37fe6694bd..5d21b46c9fda1 100644 --- a/sycl/include/sycl/kernel_bundle_enums.hpp +++ b/sycl/include/sycl/kernel_bundle_enums.hpp @@ -11,7 +11,18 @@ namespace sycl { inline namespace _V1 { -enum class bundle_state : char { input = 0, object = 1, executable = 2 }; +enum class bundle_state : char { + input = 0, + object = 1, + executable = 2, + ext_oneapi_source = 3 +}; -} +namespace ext::oneapi::experimental { + +enum class source_language : int { opencl = 0 /* sycl , spir-v, cuda */ }; + +} // namespace ext::oneapi::experimental + +} // namespace _V1 } // namespace sycl diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index e25f4536c22f3..564bad58c18d3 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -191,6 +191,7 @@ set(SYCL_SOURCES "detail/image_impl.cpp" "detail/jit_compiler.cpp" "detail/jit_device_binaries.cpp" + "detail/kernel_compiler/kernel_compiler_opencl.cpp" "detail/kernel_impl.cpp" "detail/kernel_program_cache.cpp" "detail/memory_manager.cpp" diff --git a/sycl/source/detail/common.cpp b/sycl/source/detail/common.cpp index 6cef67e8c924c..f028b7cfbf31c 100644 --- a/sycl/source/detail/common.cpp +++ b/sycl/source/detail/common.cpp @@ -74,18 +74,14 @@ const char *stringifyErrorCode(pi_int32 error) { std::vector split_string(const std::string &str, char delimeter) { std::vector result; size_t beg = 0; - size_t length = 0; - for (const auto &x : str) { - if (x == delimeter) { - result.push_back(str.substr(beg, length)); - beg += length + 1; - length = 0; - continue; - } - length++; + size_t end = 0; + while ((end = str.find(delimeter, beg)) != std::string::npos) { + result.push_back(str.substr(beg, end - beg)); + beg = end + 1; } - if (length != 0) { - result.push_back(str.substr(beg, length)); + end = str.find('\0'); + if (beg < end) { + result.push_back(str.substr(beg, end - beg)); } return result; } diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 6f28793b6d999..109e08695c80b 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -9,11 +9,13 @@ #pragma once #include +#include #include #include #include #include #include +#include #include #include #include @@ -42,7 +44,11 @@ static bool checkAllDevicesHaveAspect(const std::vector &Devices, [&Aspect](const device &Dev) { return Dev.has(Aspect); }); } -// The class is an impl counterpart of the sycl::kernel_bundle. +namespace syclex = sycl::ext::oneapi::experimental; + +class kernel_impl; + +/// The class is an impl counterpart of the sycl::kernel_bundle. // It provides an access and utilities to manage set of sycl::device_images // objects. class kernel_bundle_impl { @@ -142,9 +148,10 @@ class kernel_bundle_impl { DeviceImage, MDevices, PropList)); break; case bundle_state::input: - throw sycl::runtime_error( - "Internal error. The target state should not be input", - PI_ERROR_INVALID_OPERATION); + case bundle_state::ext_oneapi_source: + throw sycl::runtime_error("Internal error. The target state should not " + "be input or ext_oneapi_source", + PI_ERROR_INVALID_OPERATION); break; } } @@ -320,6 +327,117 @@ class kernel_bundle_impl { } } + // oneapi_ext_kernel_compiler + // construct from source string + kernel_bundle_impl(const context &Context, syclex::source_language Lang, + const std::string &Src) + : MContext(Context), MDevices(Context.get_devices()), + MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {} + + // oneapi_ext_kernel_compiler + // interop constructor + kernel_bundle_impl(context Ctx, std::vector Devs, + device_image_plain &DevImage, + std::vector KNames) + : kernel_bundle_impl(Ctx, Devs, DevImage) { + MState = bundle_state::executable; + KernelNames = KNames; + } + + std::shared_ptr + build_from_source(const std::vector Devices, + const std::vector &BuildOptions, + std::string *LogPtr) { + assert(MState == bundle_state::ext_oneapi_source && + "bundle_state::ext_oneapi_source required"); + assert(Language == syclex::source_language::opencl && + "TODO: add other Languages. Must be OpenCL"); + + // if successful, the log is empty. if failed, throws an error with the + // compilation log. + auto spirv = + syclex::detail::OpenCLC_to_SPIRV(this->Source, BuildOptions, LogPtr); + + // see also program_manager.cpp::createSpirvProgram() + using ContextImplPtr = std::shared_ptr; + sycl::detail::pi::PiProgram PiProgram = nullptr; + ContextImplPtr ContextImpl = getSyclObjImpl(MContext); + const PluginPtr &Plugin = ContextImpl->getPlugin(); + Plugin->call( + ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &PiProgram); + + Plugin->call(PiProgram); + + for (const auto &SyclDev : Devices) { + pi::PiDevice Dev = getSyclObjImpl(SyclDev)->getHandleRef(); + Plugin->call( + PiProgram, 1, &Dev, nullptr, nullptr, nullptr); + } + + // Get the number of kernels in the program. + size_t NumKernels; + Plugin->call( + PiProgram, PI_PROGRAM_INFO_NUM_KERNELS, sizeof(size_t), &NumKernels, + nullptr); + + // Get the kernel names. + size_t KernelNamesSize; + Plugin->call( + PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, 0, nullptr, &KernelNamesSize); + + // semi-colon delimited list of kernel names. + std::string KernelNamesStr(KernelNamesSize, ' '); + Plugin->call( + PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, KernelNamesStr.size(), + &KernelNamesStr[0], nullptr); + std::vector KernelNames = + detail::split_string(KernelNamesStr, ';'); + + // make the device image and the kernel_bundle_impl + auto KernelIDs = std::make_shared>(); + auto DevImgImpl = std::make_shared( + nullptr, MContext, MDevices, bundle_state::executable, KernelIDs, + PiProgram); + device_image_plain DevImg{DevImgImpl}; + return std::make_shared(MContext, MDevices, DevImg, + KernelNames); + } + + bool ext_oneapi_has_kernel(const std::string &Name) { + auto it = std::find(KernelNames.begin(), KernelNames.end(), Name); + return it != KernelNames.end(); + } + + kernel + ext_oneapi_get_kernel(const std::string &Name, + const std::shared_ptr &Self) { + if (KernelNames.empty()) + throw sycl::exception(make_error_code(errc::invalid), + "'ext_oneapi_get_kernel' is only available in " + "kernel_bundles successfully built from " + "kernel_bundle."); + + if (!ext_oneapi_has_kernel(Name)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + Name + "' not found in kernel_bundle"); + + assert(MDeviceImages.size() > 0); + const std::shared_ptr &DeviceImageImpl = + detail::getSyclObjImpl(MDeviceImages[0]); + sycl::detail::pi::PiProgram PiProgram = DeviceImageImpl->get_program_ref(); + ContextImplPtr ContextImpl = getSyclObjImpl(MContext); + const PluginPtr &Plugin = ContextImpl->getPlugin(); + sycl::detail::pi::PiKernel PiKernel = nullptr; + Plugin->call(PiProgram, Name.c_str(), &PiKernel); + + Plugin->call(PiKernel); + + std::shared_ptr KernelImpl = std::make_shared( + PiKernel, detail::getSyclObjImpl(MContext), Self); + + return detail::createSyclObjFromImpl(KernelImpl); + } + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { @@ -557,6 +675,11 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; + // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames + const syclex::source_language Language = syclex::source_language::opencl; + const std::string Source; + // only kernel_bundles created from source have KernelNames member. + std::vector KernelNames; }; } // namespace detail diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp new file mode 100644 index 0000000000000..bf308229a6cb8 --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -0,0 +1,196 @@ +//==-- kernel_compiler_opencl.cpp OpenCL kernel compilation support -==// +// +// 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 // getOsLibraryFuncAddress +#include // for make_error_code + +#include "kernel_compiler_opencl.hpp" + +#include "../online_compiler/ocloc_api.h" + +#include // strlen +#include // for std::accumulate + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +// ensures the OclocLibrary has the right version, etc. +void checkOclocLibrary(void *OclocLibrary) { + void *OclocVersionHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion"); + // The initial versions of ocloc library did not have the oclocVersion() + // function. Those versions had the same API as the first version of ocloc + // library having that oclocVersion() function. + int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0; + if (OclocVersionHandle) { + decltype(::oclocVersion) *OclocVersionFunc = + reinterpret_cast(OclocVersionHandle); + LoadedVersion = OclocVersionFunc(); + } + // The loaded library with version (A.B) is compatible with expected API/ABI + // version (X.Y) used here if A == B and B >= Y. + int LoadedVersionMajor = LoadedVersion >> 16; + int LoadedVersionMinor = LoadedVersion & 0xffff; + int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16; + int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff; + if (LoadedVersionMajor != CurrentVersionMajor || + LoadedVersionMinor < CurrentVersionMinor) { + throw sycl::exception( + make_error_code(errc::build), + std::string("Found incompatible version of ocloc library: (") + + std::to_string(LoadedVersionMajor) + "." + + std::to_string(LoadedVersionMinor) + + "). The supported versions are (" + + std::to_string(CurrentVersionMajor) + + ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ")."); + } +} + +static void *OclocLibrary = nullptr; + +// load the ocloc shared library, check it. +void *loadOclocLibrary() { +#ifdef __SYCL_RT_OS_WINDOWS + static const std::string OclocLibraryName = "ocloc64.dll"; +#else + static const std::string OclocLibraryName = "libocloc.so"; +#endif + void *tempPtr = OclocLibrary; + if (tempPtr == nullptr) { + tempPtr = sycl::detail::pi::loadOsLibrary(OclocLibraryName); + + if (tempPtr == nullptr) + throw sycl::exception(make_error_code(errc::build), + "Unable to load ocloc library " + OclocLibraryName); + + checkOclocLibrary(tempPtr); + + OclocLibrary = tempPtr; + } + + return OclocLibrary; +} + +bool OpenCLC_Compilation_Available() { + // Already loaded? + if (OclocLibrary != nullptr) + return true; + + try { + // loads and checks version + loadOclocLibrary(); + return true; + } catch (...) { + return false; + } +} + +spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, + const std::vector &UserArgs, + std::string *LogPtr) { + std::vector CMUserArgs = UserArgs; + CMUserArgs.push_back("-cmc"); + + // handles into ocloc shared lib + static void *oclocInvokeHandle = nullptr; + static void *oclocFreeOutputHandle = nullptr; + std::error_code build_errc = make_error_code(errc::build); + + // setup Library + if (!oclocInvokeHandle) { + if (OclocLibrary == nullptr) + loadOclocLibrary(); + + oclocInvokeHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); + if (!oclocInvokeHandle) + throw sycl::exception(build_errc, "Cannot load oclocInvoke() function"); + + oclocFreeOutputHandle = sycl::detail::pi::getOsLibraryFuncAddress( + OclocLibrary, "oclocFreeOutput"); + if (!oclocFreeOutputHandle) + throw sycl::exception(build_errc, + "Cannot load oclocFreeOutput() function"); + } + + // assemble ocloc args + std::string CombinedUserArgs = + std::accumulate(UserArgs.begin(), UserArgs.end(), std::string(""), + [](const std::string &acc, const std::string &s) { + return acc + s + " "; + }); + + std::vector Args = {"ocloc", "-q", "-spv_only", "-options", + CombinedUserArgs.c_str()}; + + uint32_t NumOutputs = 0; + uint8_t **Outputs = nullptr; + uint64_t *OutputLengths = nullptr; + char **OutputNames = nullptr; + + const uint8_t *Sources[] = { + reinterpret_cast(Source.c_str())}; + const char *SourceName = "main.cl"; + const uint64_t SourceLengths[] = {Source.length() + 1}; + + Args.push_back("-file"); + Args.push_back(SourceName); + + // invoke + decltype(::oclocInvoke) *OclocInvokeFunc = + reinterpret_cast(oclocInvokeHandle); + int CompileError = + OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths, + &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs, + &Outputs, &OutputLengths, &OutputNames); + + // gather the results ( the SpirV and the Log) + spirv_vec_t SpirV; + std::string CompileLog; + for (uint32_t i = 0; i < NumOutputs; i++) { + size_t NameLen = strlen(OutputNames[i]); + if (NameLen >= 4 && strstr(OutputNames[i], ".spv") != nullptr && + Outputs[i] != nullptr) { + assert(SpirV.size() == 0 && "More than one SPIR-V output found."); + SpirV = spirv_vec_t(Outputs[i], Outputs[i] + OutputLengths[i]); + } else if (!strcmp(OutputNames[i], "stdout.log")) { + const char *LogText = reinterpret_cast(Outputs[i]); + if (LogText != nullptr && LogText[0] != '\0') { + CompileLog.append(LogText); + if (LogPtr != nullptr) + LogPtr->append(LogText); + } + } + } + + // Try to free memory before reporting possible error. + decltype(::oclocFreeOutput) *OclocFreeOutputFunc = + reinterpret_cast(oclocFreeOutputHandle); + int MemFreeError = + OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames); + + if (CompileError) + throw sycl::exception(build_errc, "ocloc reported compilation errors: {\n" + + CompileLog + "\n}"); + + if (SpirV.empty()) + throw sycl::exception(build_errc, + "Unexpected output: ocloc did not return SPIR-V"); + + if (MemFreeError) + throw sycl::exception(build_errc, "ocloc cannot safely free resources"); + + return SpirV; +} + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp new file mode 100644 index 0000000000000..8826cf5ea27a7 --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -0,0 +1,34 @@ +//==-- kernel_compiler_opencl.hpp OpenCL kernel compilation support -==// +// +// 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 // for __SYCL_EXPORT +#include + +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +using spirv_vec_t = std::vector; +spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, + const std::vector &UserArgs, + std::string *LogPtr); + +bool OpenCLC_Compilation_Available(); + +} // namespace detail +} // namespace ext::oneapi::experimental + +} // 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 de4bfb86e8da0..9d66d71a713ca 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1910,7 +1910,15 @@ void ProgramManager::bringSYCLDeviceImagesToState( for (device_image_plain &DevImage : DeviceImages) { const bundle_state DevImageState = getSyclObjImpl(DevImage)->get_state(); + // At this time, there is no circumstance where a device image should ever + // be in the source state. That not good. + assert(DevImageState != bundle_state::ext_oneapi_source); + switch (TargetState) { + case bundle_state::ext_oneapi_source: + // This case added for switch statement completion. We should not be here. + assert(DevImageState == bundle_state::ext_oneapi_source); + break; case bundle_state::input: // Do nothing since there is no state which can be upgraded to the input. assert(DevImageState == bundle_state::input); @@ -1926,6 +1934,11 @@ void ProgramManager::bringSYCLDeviceImagesToState( break; case bundle_state::executable: { switch (DevImageState) { + case bundle_state::ext_oneapi_source: + // This case added for switch statement completion. + // We should not be here. + assert(DevImageState != bundle_state::ext_oneapi_source); + break; case bundle_state::input: DevImage = build(DevImage, getSyclObjImpl(DevImage)->get_devices(), /*PropList=*/{}); diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index dd461b7049aa7..e10272cf391c7 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -43,6 +43,8 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_BINDLESS_IMAGES 1 #define SYCL_EXT_ONEAPI_GROUP_ALGORITHMS 1 #define SYCL_EXT_ONEAPI_GROUP_SORT 1 +#define SYCL_EXT_ONEAPI_KERNEL_COMPILER 1 +#define SYCL_EXT_ONEAPI_KERNEL_COMPILER_OPENCL 1 #define SYCL_EXT_ONEAPI_MAX_WORK_GROUP_QUERY 1 #define SYCL_EXT_ONEAPI_ND_RANGE_REDUCTIONS 1 #define SYCL_EXT_ONEAPI_DEFAULT_CONTEXT 1 diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 350bbd4260ff5..1072719be9653 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -217,6 +217,7 @@ event handler::finalize() { // Nothing to do break; case bundle_state::object: + case bundle_state::ext_oneapi_source: assert(0 && "Expected that the bundle is either in input or executable " "states."); break; diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e656ec555113b..4b8e2551d4517 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -8,6 +8,7 @@ #include #include +#include #include #include @@ -111,6 +112,14 @@ bool kernel_bundle_plain::is_specialization_constant_set( return impl->is_specialization_constant_set(SpecName); } +bool kernel_bundle_plain::ext_oneapi_has_kernel(const std::string &name) { + return impl->ext_oneapi_has_kernel(name); +} + +kernel kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) { + return impl->ext_oneapi_get_kernel(name, impl); +} + ////////////////////////////////// ///// sycl::detail free functions ////////////////////////////////// @@ -346,5 +355,69 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { return true; } +///////////////////////// +// * kernel_compiler extension * +///////////////////////// +namespace ext::oneapi::experimental { + +using source_kb = kernel_bundle; +using exe_kb = kernel_bundle; +using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; + +///////////////////////// +// syclex::is_source_kernel_bundle_supported +///////////////////////// +bool is_source_kernel_bundle_supported(backend BE, source_language Language) { + // at the moment, OpenCL is the only language supported + // and it's 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 ((Language == source_language::opencl) && BE_Acceptable) { + return detail::OpenCLC_Compilation_Available(); + } + + // otherwise + return false; +} + +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +source_kb create_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::string &Source) { + // TODO: if we later support a "reason" why support isn't present + // (like a missing shared library etc.) it'd be nice to include it in + // the exception message here. + backend BE = SyclContext.get_backend(); + if (!is_source_kernel_bundle_supported(BE, Language)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel_bundle creation from source not supported"); + + std::shared_ptr KBImpl = + std::make_shared(SyclContext, Language, Source); + return sycl::detail::createSyclObjFromImpl(KBImpl); +} + +///////////////////////// +// syclex::detail::build_from_source(source_kb) => exe_kb +///////////////////////// +namespace detail { + +exe_kb build_from_source(source_kb &SourceKB, + const std::vector &Devices, + const std::vector &BuildOptions, + std::string *LogPtr) { + std::vector UniqueDevices = + sycl::detail::removeDuplicateDevices(Devices); + std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); + std::shared_ptr KBImpl = + sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr); + return sycl::detail::createSyclObjFromImpl(KBImpl); +} + +} // namespace detail +} // namespace ext::oneapi::experimental + } // namespace _V1 } // namespace sycl diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp new file mode 100644 index 0000000000000..9bfce76e06eb9 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -0,0 +1,167 @@ +//==- kernel_compiler.cpp --- kernel_compiler extension tests -------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: ocloc + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// the new spec for the kernel_compiler opens the door to supporting several +// different source languages. But, initially, OpenCL Kernels are the only ones +// supported. This test is limited to that (thus the cm-compiler requirement) +// but in the future it may need to broken out into other tests. + +#include + +auto constexpr CLSource = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0); + out[i] = in[i]*2 + 100; +} +__kernel void her_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0); + out[i] = in[i]*5 + 1000; +} +)==="; + +auto constexpr BadCLSource = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0) + no semi-colon!! + out[i] = in[i]*2 + 100; +} +)==="; +/* +Compile Log: +1:3:34: error: use of undeclared identifier 'no' + size_t i = get_global_id(0) + no semi-colon!! + ^ +1:3:36: error: expected ';' at end of declaration + size_t i = get_global_id(0) + no semi-colon!! + ^ + ; + +Build failed with error code: -11 + +============= + +*/ + +using namespace sycl; + +void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, + int added) { + constexpr int N = 4; + cl_int InputArray[N] = {0, 1, 2, 3}; + cl_int OutputArray[N] = {}; + + sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); + sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); + + Q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }); + + sycl::host_accessor Out{OutputBuf}; + for (int I = 0; I < N; I++) + assert(Out[I] == ((I * multiplier) + added)); +} + +void test_build_and_run() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = syclex::is_source_kernel_bundle_supported( + ctx.get_backend(), syclex::source_language::opencl); + if (!ok) { + std::cout << "Apparently this backend does not support OpenCL C source " + "kernel bundle extension: " + << ctx.get_backend() << std::endl; + return; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, CLSource); + // compilation of empty prop list, no devices + exe_kb kbExe1 = syclex::build(kbSrc); + + // compilation with props and devices + std::string log; + std::vector flags{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}; + std::vector devs = kbSrc.get_devices(); + sycl::context ctxRes = kbSrc.get_context(); + assert(ctxRes == ctx); + sycl::backend beRes = kbSrc.get_backend(); + assert(beRes == ctx.get_backend()); + + exe_kb kbExe2 = syclex::build( + kbSrc, devs, + syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); + + bool hasMyKernel = kbExe2.ext_oneapi_has_kernel("my_kernel"); + bool hasHerKernel = kbExe2.ext_oneapi_has_kernel("her_kernel"); + bool notExistKernel = kbExe2.ext_oneapi_has_kernel("not_exist"); + assert(hasMyKernel && "my_kernel should exist, but doesn't"); + assert(hasHerKernel && "her_kernel should exist, but doesn't"); + assert(!notExistKernel && "non-existing kernel should NOT exist, but does?"); + + sycl::kernel my_kernel = kbExe2.ext_oneapi_get_kernel("my_kernel"); + sycl::kernel her_kernel = kbExe2.ext_oneapi_get_kernel("her_kernel"); + + auto my_num_args = my_kernel.get_info(); + assert(my_num_args == 2 && "my_kernel should take 2 args"); + + testSyclKernel(q, my_kernel, 2, 100); + testSyclKernel(q, her_kernel, 5, 1000); +} + +void test_error() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = syclex::is_source_kernel_bundle_supported( + ctx.get_backend(), syclex::source_language::opencl); + if (!ok) { + return; + } + + try { + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, BadCLSource); + exe_kb kbExe1 = syclex::build(kbSrc); + assert(false && "we should not be here."); + } catch (sycl::exception &e) { + // nice! + assert(e.code() == sycl::errc::build); + } + // any other error will escape and cause the test to fail ( as it should ). +} + +int main() { +#ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_OPENCL + static_assert(false, "KernelCompiler OpenCL feature test macro undefined"); +#endif + +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + test_build_and_run(); + test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp new file mode 100644 index 0000000000000..65dcbe6eddb0a --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -0,0 +1,181 @@ +//==- opencl_capabilities.cpp ----------------------------------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: ocloc + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Here we are testing some of the various args that SYCL can and cannot +// pass to an OpenCL kernel that is compiled with the kernel_compiler. + +// IMPORTANT: LevelZero YES! +// Even though this test is covering which OpenCL capabilities +// are covered by the kernel_compiler, this is not a test of only +// the OpenCL devices. The LevelZero backend works with the kernel_compiler +// so long as ocloc is installed and should be able to +// successfully run and pass these tests. + +#include +using namespace sycl; +namespace syclex = sycl::ext::oneapi::experimental; +using source_kb = sycl::kernel_bundle; +using exe_kb = sycl::kernel_bundle; + +// ----------------------- +// local accessor +// ----------------------- +auto constexpr LocalAccCLSource = R"===( + kernel void test_la(global int *buf, local float *slm, int n) { + if (get_local_id(0) == 0) { + for (int i = 0; i < n; i++) + slm[i] = i + get_group_id(0); + } + barrier(CLK_LOCAL_MEM_FENCE); + + bool ok = true; + for (int i = 0; i < n; i++) + ok &= (slm[i] == i + get_group_id(0)); + + buf[get_global_id(0)] = ok; + } +)==="; + +void test_local_accessor() { + + sycl::queue q; + sycl::context ctx = q.get_context(); + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, LocalAccCLSource); + exe_kb kbExe1 = syclex::build(kbSrc); + sycl::kernel test_kernel = kbExe1.ext_oneapi_get_kernel("test_la"); + + constexpr cl_int N_slm = 256; + constexpr int N_wg = 32; + + cl_int init[N_wg]; + sycl::buffer b(init, N_wg); + + q.submit([&](handler &cgh) { + auto acc_global = b.get_access(cgh); + local_accessor acc_local(N_slm, cgh); + + cgh.set_arg(0, acc_global); + cgh.set_arg(1, acc_local); + cgh.set_arg(2, N_slm); + + cgh.parallel_for(nd_range<1>(N_wg, 1), test_kernel); + }); + + sycl::host_accessor Out{b}; + for (int i = 0; i < N_wg; i++) + assert(Out[i] == 1); +} + +// ----------------------- +// USM pointer and scalars +// ----------------------- +auto constexpr USMCLSource = R"===( +__kernel void usm_kernel(__global int *usmPtr, int multiplier, float added) { + size_t i = get_global_id(0); + usmPtr[i] = (i * multiplier) + added; +} +)==="; + +void test_usm_pointer_and_scalar() { + sycl::queue q; + sycl::context ctx = q.get_context(); + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, USMCLSource); + exe_kb kbExe1 = syclex::build(kbSrc); + sycl::kernel usm_kernel = kbExe1.ext_oneapi_get_kernel("usm_kernel"); + + // the scalars submitted to the kernel + cl_int multiplier = 2; + cl_float added = 100.f; + constexpr size_t N = 32; + cl_int *usmPtr = sycl::malloc_shared(N, q); + + q.submit([&](sycl::handler &cgh) { + cgh.set_arg(0, usmPtr); + cgh.set_arg(1, multiplier); // scalar args + cgh.set_arg(2, added); + cgh.parallel_for(sycl::range<1>{N}, usm_kernel); + }); + q.wait(); + + for (size_t i = 0; i < N; i++) { + assert(usmPtr[i] == ((i * multiplier) + added)); + } + + sycl::free(usmPtr, ctx); +} + +// ----------------------- +// structure passed by value +// Note that it is imperative that the struct defined in the OpenCL C string +// exactly match the one used for any kernel args. Overlooking their duality +// will lead to difficult to discover errors. +// ----------------------- + +auto constexpr StructSrc = R"===( +struct pair { + int multiplier; + float added; +}; +__kernel void struct_kernel(__global int *usmPtr, struct pair adjuster) { + size_t i = get_global_id(0); + usmPtr[i] = (i * adjuster.multiplier) + adjuster.added; +} +)==="; + +struct pair { + cl_int multiplier; + cl_float added; +}; + +void test_struct() { + sycl::queue q; + sycl::context ctx = q.get_context(); + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, StructSrc); + exe_kb kbExe1 = syclex::build(kbSrc); + sycl::kernel struct_kernel = kbExe1.ext_oneapi_get_kernel("struct_kernel"); + + pair adjuster; + adjuster.multiplier = 2, adjuster.added = 100.f; + constexpr size_t N = 32; + cl_int *usmPtr = sycl::malloc_shared(N, q); + + q.submit([&](sycl::handler &cgh) { + cgh.set_arg(0, usmPtr); + cgh.set_arg(1, adjuster); // struct by value + cgh.parallel_for(sycl::range<1>{N}, struct_kernel); + }); + q.wait(); + + for (size_t i = 0; i < N; i++) { + assert(usmPtr[i] == ((i * adjuster.multiplier) + adjuster.added)); + } + + sycl::free(usmPtr, ctx); +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER_OPENCL + test_local_accessor(); + test_usm_pointer_and_scalar(); + test_struct(); +#else + static_assert(false, "KernelCompiler OpenCL feature test macro undefined"); +#endif + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 49c71643e7a86..56cd6ebf90be1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3713,10 +3713,13 @@ _ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_m _ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_mem_handleERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental32create_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE +_ZN4sycl3_V13ext6oneapi12experimental33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC2ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD1Ev _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD2Ev +_ZN4sycl3_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bundleILNS0_12bundle_stateE3EEERKSt6vectorINS0_6deviceESaISA_EERKS9_INSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISK_EEPSK_ _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph12finalizeImplEv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graphC1ERKSt10shared_ptrINS4_10graph_implEERKNS0_7contextE @@ -3981,6 +3984,8 @@ _ZN4sycl3_V16detail19convertChannelOrderE23_pi_image_channel_order _ZN4sycl3_V16detail19convertChannelOrderENS0_19image_channel_orderE _ZN4sycl3_V16detail19getImageElementSizeEhNS0_18image_channel_typeE _ZN4sycl3_V16detail19getPluginOpaqueDataILNS0_7backendE5EEEPvS4_ +_ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_get_kernelERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE +_ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_has_kernelERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN4sycl3_V16detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_16AccessorBaseHostENS0_6access6targetE _ZN4sycl3_V16detail20associateWithHandlerERNS0_7handlerEPNS1_28SampledImageAccessorBaseHostENS0_12image_targetE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 63b57152d2127..ecc39d00ff59c 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -120,8 +120,8 @@ ??$get_info@Ukernel_kernel_pipe_support@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ ??$get_info@Ulocal_mem_size@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Ulocal_mem_type@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA?AW4local_mem_type@info@12@XZ -??$get_info@Umatrix_combinations@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@Umatrix_combinations@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$vector@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info@Umatrix_combinations@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Ucombination@matrix@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@Umax_clock_frequency@device@info@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ ??$get_info@Umax_compute_queue_indices@device@info@intel@ext@_V1@sycl@@@device@_V1@sycl@@QEBAHXZ ??$get_info@Umax_compute_queue_indices@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAHXZ @@ -932,6 +932,7 @@ ?begin@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA_NAEAVqueue@67@@Z ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA_NAEBV?$vector@Vqueue@_V1@sycl@@V?$allocator@Vqueue@_V1@sycl@@@std@@@std@@@Z +?build_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$01@56@AEAV?$kernel_bundle@$02@56@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@PEAV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?build_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$kernel_bundle@$0A@@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z ?canReadHostPtr@SYCLMemObjT@detail@_V1@sycl@@QEAA_NPEAX_K@Z ?canReuseHostPtr@SYCLMemObjT@detail@_V1@sycl@@QEAA_NPEAX_K@Z @@ -974,6 +975,7 @@ ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@AEAVimage_mem@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z ?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?create_kernel_bundle_from_source@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@45@AEBVcontext@45@W4source_language@12345@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?default_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z ?deleteAccProps@buffer_plain@detail@_V1@sycl@@IEAAXAEBW4PropWithDataKind@234@@Z ?deleteAccessorProperty@SYCLMemObjT@detail@_V1@sycl@@QEAAXAEBW4PropWithDataKind@234@@Z @@ -1039,11 +1041,13 @@ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ +?ext_oneapi_get_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA?AVkernel@34@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_has_kernel@kernel_bundle_plain@detail@_V1@sycl@@QEAA_NAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?ext_oneapi_memcpy2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_memset2d_impl@handler@_V1@sycl@@AEAAXPEAX_KH11@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vcontext@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vcontext@_V1@sycl@@@2oneapi@ext@34@@Z @@ -1335,6 +1339,7 @@ ?is_host@queue@_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@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@45@W4source_language@12345@@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 ?link_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$kernel_bundle@$00@_V1@sycl@@V?$allocator@V?$kernel_bundle@$00@_V1@sycl@@@std@@@5@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp new file mode 100644 index 0000000000000..20f58009ee968 --- /dev/null +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -0,0 +1,109 @@ +//==- kernel_compiler_constraints.cpp --------------------------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// RUN: %clangxx -fsyntax-only -fsycl -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +// kernel_bundles sporting the new bundle_state::ext_oneapi_source should NOT +// support several member functions. This test confirms that. + +#include + +int main() { +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + std::vector devices = ctx.get_devices(); + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, ""); + + // expected-error@+1 {{no matching member function for call to 'contains_specialization_constants'}} + kbSrc.contains_specialization_constants(); + + // expected-error@+1 {{no matching member function for call to 'native_specialization_constant'}} + kbSrc.native_specialization_constant(); + + constexpr sycl::specialization_id SpecName; + // expected-error@+1 {{no matching member function for call to 'has_specialization_constant'}} + kbSrc.has_specialization_constant(); + + // expected-error@+1 {{no matching member function for call to 'get_specialization_constant'}} + auto i = kbSrc.get_specialization_constant(); + + // expected-error@+1 {{no matching member function for call to 'get_kernel'}} + kbSrc.get_kernel(); + + // expected-error@+1 {{no matching member function for call to 'get_kernel_ids'}} + std::vector vec = kbSrc.get_kernel_ids(); + + class TestKernel1; + sycl::kernel_id TestKernel1ID = sycl::get_kernel_id(); + + // expected-error@+1 {{no matching member function for call to 'has_kernel'}} + kbSrc.has_kernel(); + + // expected-error@+1 {{no matching member function for call to 'has_kernel'}} + kbSrc.has_kernel(devices[0]); + + // expected-error@+1 {{no matching member function for call to 'has_kernel'}} + kbSrc.has_kernel(TestKernel1ID); + + // expected-error@+1 {{no matching member function for call to 'has_kernel'}} + kbSrc.has_kernel(TestKernel1ID, devices[0]); + + // expected-error@+1 {{no matching member function for call to 'begin'}} + kbSrc.begin(); + + // expected-error@+1 {{no matching member function for call to 'end'}} + kbSrc.end(); + + // expected-error@+1 {{no matching member function for call to 'empty'}} + kbSrc.empty(); + + std::string log; + std::vector flags{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}; + // OK + syclex::build(kbSrc); + + // expected-error@+1 {{no matching function for call to 'build'}} + syclex::build(kbSrc, + syclex::properties{syclex::usm_kind}); + + // OK + syclex::build(kbSrc, syclex::properties{syclex::build_options{flags}, + syclex::save_log{&log}}); + + // expected-error@../include/sycl/ext/oneapi/properties/properties.hpp:* {{too many template arguments for class template 'all_props_are_keys_of'}} + // expected-error@+1 {{no matching function for call to 'build'}} + syclex::build(kbSrc, syclex::properties{ + syclex::build_options{flags}, syclex::save_log{&log}, + syclex::usm_kind}); + // OK + syclex::build(kbSrc, devices); + + // expected-error@+1 {{no matching function for call to 'build'}} + syclex::build(kbSrc, devices, + syclex::properties{syclex::usm_kind}); + + // OK + syclex::build( + kbSrc, devices, + syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); + + // expected-error@+1 {{no matching function for call to 'build'}} + syclex::build(kbSrc, devices, + syclex::properties{syclex::build_options{flags}, + syclex::save_log{&log}, + syclex::usm_kind}); + +#endif +} diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp index 0d8cef657755d..eefee088edd05 100644 --- a/sycl/test/warnings/sycl_2020_deprecations.cpp +++ b/sycl/test/warnings/sycl_2020_deprecations.cpp @@ -118,6 +118,11 @@ int main() { ex.get_cl_code(); (void)ex; + // expected-warning@+1{{'online_compiler' is deprecated}} + sycl::ext::intel::experimental::online_compiler< + sycl::ext::intel::experimental::source_language::opencl_c> + oc(Device); + Queue.submit([](sycl::handler &CGH) { // expected-warning@+3{{'nd_range' is deprecated: offsets are deprecated in SYCL2020}} // expected-warning@+2{{'nd_range' is deprecated: offsets are deprecated in SYCL2020}}