From 543848c74e30d61ae7add15d628e7f3ded354e62 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 26 Sep 2023 12:11:06 -0700 Subject: [PATCH 01/37] who doesn't love enums? --- sycl/include/sycl/kernel_bundle.hpp | 10 ++++++++++ sycl/include/sycl/kernel_bundle_enums.hpp | 15 +++++++++++++-- sycl/source/detail/kernel_bundle_impl.hpp | 7 ++++--- .../detail/program_manager/program_manager.cpp | 13 +++++++++++++ sycl/source/handler.cpp | 1 + 5 files changed, 41 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index fdd2b0c77bd1b..4b167ae8a33fb 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -743,6 +743,16 @@ build(const kernel_bundle &InputBundle, return build(InputBundle, InputBundle.get_devices(), PropList); } +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +namespace ext::oneapi::experimental { +namespace syclex = ext::oneapi::experimental; +kernel_bundle +create_kernel_bundle_from_source(const context &syclContext, + syclex::source_language, std::string &source); +} // 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/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index a990a1efe579b..7c9ddd11f78fc 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -142,9 +142,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; } } diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 4950c6539bb41..8d69193f69404 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1893,7 +1893,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); @@ -1909,6 +1917,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/handler.cpp b/sycl/source/handler.cpp index 83d2250c970cf..560d8fd019458 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; From 01644683aa90053d375b0f874975dbf6c79db5a6 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 26 Sep 2023 16:49:01 -0700 Subject: [PATCH 02/37] create_kernel_bundle_from_source() --- sycl/include/sycl/kernel_bundle.hpp | 9 +++++---- sycl/source/detail/kernel_bundle_impl.hpp | 12 ++++++++++++ sycl/source/kernel_bundle.cpp | 20 ++++++++++++++++++++ 3 files changed, 37 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 4b167ae8a33fb..ce8cdcd1d4108 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -747,10 +747,11 @@ build(const kernel_bundle &InputBundle, // syclex::create_kernel_bundle_from_source ///////////////////////// namespace ext::oneapi::experimental { -namespace syclex = ext::oneapi::experimental; -kernel_bundle -create_kernel_bundle_from_source(const context &syclContext, - syclex::source_language, std::string &source); +namespace syclex = sycl::ext::oneapi::experimental; +__SYCL_EXPORT kernel_bundle +create_kernel_bundle_from_source(const context &SyclContext, + const syclex::source_language Language, + const std::string &Source); } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 7c9ddd11f78fc..59d6a35588503 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -42,6 +42,8 @@ static bool checkAllDevicesHaveAspect(const std::vector &Devices, [&Aspect](const device &Dev) { return Dev.has(Aspect); }); } +namespace syclex = sycl::ext::oneapi::experimental; + // 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. @@ -321,6 +323,13 @@ class kernel_bundle_impl { } } + // oneapi ext kernel_compiler + // construct from source string + kernel_bundle_impl(const context &Context, const syclex::source_language Lang, + const std::string &Src) + : MContext(Context), MDevices(Context.get_devices()), + MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {} + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { @@ -519,6 +528,9 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; + // ext_oneapi_kernel_compiler : Source and Languauge + const syclex::source_language Language = syclex::source_language::opencl; + const std::string Source; }; } // namespace detail diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e656ec555113b..e5b2f712eec03 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -346,5 +346,25 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { return true; } +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +namespace ext::oneapi::experimental { + +namespace syclex = sycl::ext::oneapi::experimental; +using source_kb = kernel_bundle; +using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; + +source_kb +create_kernel_bundle_from_source(const context &SyclContext, + const syclex::source_language Language, + const std::string &Source) { + std::shared_ptr KBImpl = + std::make_shared(SyclContext, Language, Source); + return sycl::detail::createSyclObjFromImpl(KBImpl); +} + +} // namespace ext::oneapi::experimental + } // namespace _V1 } // namespace sycl From fc8b8978044c2cb7df3fc590013b9ab349f98579 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 27 Sep 2023 10:25:36 -0700 Subject: [PATCH 03/37] is_source_kernel_bundle_supported plus interim --- sycl/include/sycl/kernel_bundle.hpp | 20 +++++++-- sycl/source/detail/kernel_bundle_impl.hpp | 20 ++++++++- sycl/source/kernel_bundle.cpp | 49 ++++++++++++++++++++--- 3 files changed, 79 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index ce8cdcd1d4108..84257b76eba72 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -743,15 +743,29 @@ build(const kernel_bundle &InputBundle, return build(InputBundle, InputBundle.get_devices(), PropList); } +namespace ext::oneapi::experimental { +namespace syclex = sycl::ext::oneapi::experimental; +///////////////////////// +// 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 ///////////////////////// -namespace ext::oneapi::experimental { -namespace syclex = sycl::ext::oneapi::experimental; __SYCL_EXPORT kernel_bundle create_kernel_bundle_from_source(const context &SyclContext, - const syclex::source_language Language, + syclex::source_language Language, const std::string &Source); + +///////////////////////// +// syclex::build(source_kb) => exe_kb +///////////////////////// +__SYCL_EXPORT kernel_bundle +build(kernel_bundle &SourceKB, + const property_list &PropList = {}); + } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 59d6a35588503..f13e6af522e86 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -44,7 +44,7 @@ static bool checkAllDevicesHaveAspect(const std::vector &Devices, namespace syclex = sycl::ext::oneapi::experimental; -// The class is an impl counterpart of the sycl::kernel_bundle. +/// 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 { @@ -325,11 +325,21 @@ class kernel_bundle_impl { // oneapi ext kernel_compiler // construct from source string - kernel_bundle_impl(const context &Context, const syclex::source_language Lang, + 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 + // construct exe from source kb. + kernel_bundle_impl( + const kernel_bundle &SourceBundle) + : MContext(SourceBundle.get_context()), + MDevices(SourceBundle.get_devices()), MState(bundle_state::executable) { + + // sourceImpl = getSyclObjImpl(SourceBundle); + } + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { @@ -531,6 +541,12 @@ class kernel_bundle_impl { // ext_oneapi_kernel_compiler : Source and Languauge const syclex::source_language Language = syclex::source_language::opencl; const std::string Source; + + // friend declaration for build(source_kb) is a wee ungainly + friend kernel_bundle + sycl::ext::oneapi::experimental::build( + kernel_bundle &SourceKB, + const property_list &PropList); }; } // namespace detail diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e5b2f712eec03..a622d9c0c5c7c 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -11,6 +11,8 @@ #include #include +// #include + #include namespace sycl { @@ -347,23 +349,60 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { } ///////////////////////// -// syclex::create_kernel_bundle_from_source +// * kernel_compiler extension * ///////////////////////// namespace ext::oneapi::experimental { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = kernel_bundle; +using exe_kb = kernel_bundle; using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; -source_kb -create_kernel_bundle_from_source(const context &SyclContext, - const syclex::source_language Language, - const std::string &Source) { +///////////////////////// +// syclex::is_source_kernel_bundle_supported +///////////////////////// +bool is_source_kernel_bundle_supported(backend BE, source_language Language) { + // TODO - maybe return false? + return true; +} + +///////////////////////// +// syclex::create_kernel_bundle_from_source +///////////////////////// +source_kb create_kernel_bundle_from_source(const context &SyclContext, + syclex::source_language Language, + const std::string &Source) { + // TODO -- throw errc::invalid if lang is not supported by BE. + // use syclex::is_source_kernel_bundle_supported(BE, Lang) std::shared_ptr KBImpl = std::make_shared(SyclContext, Language, Source); return sycl::detail::createSyclObjFromImpl(KBImpl); } +///////////////////////// +// syclex::build(source_kb) => exe_kb +///////////////////////// + +exe_kb build(source_kb &SourceKB, const property_list &PropList) { + + // CP gross test code. Using existing OnlineCompiler as placeholder. + std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); + // sycl::ext::intel::experimental::online_compiler + // Compiler; + std::vector flags{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}; + std::string s = sourceImpl->Source; + std::cout << "sourcey: " << s << std::endl; + // std::vector SpirVec = Compiler.compile(sourceImpl->Source, flags); + + backend BE = SourceKB.get_backend(); + + // CP fake code to compile for the nonce. This constructor is empty. + std::shared_ptr KBImpl = + std::make_shared(SourceKB); + return sycl::detail::createSyclObjFromImpl(KBImpl); +} + } // namespace ext::oneapi::experimental } // namespace _V1 From 0759ce9236500758c18f3541766efdb7fa2fb2e7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 27 Sep 2023 15:02:22 -0700 Subject: [PATCH 04/37] OCL SpirV gen. Need templating, wrapping for language dispatch, and more. But, most importantly, the PI which is the biggest hurdle this will face and is my next target. Better get a running start. --- .../kernel_compiler_opencl.hpp | 166 ++++++++++++++++++ sycl/source/kernel_bundle.cpp | 5 +- 2 files changed, 170 insertions(+), 1 deletion(-) create mode 100644 sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp 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..b7ad41b17d19c --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -0,0 +1,166 @@ +//==-- 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 "../online_compiler/ocloc_api.h" + +#include // for std::accumulate + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +// copy/pasta from online_compiler.cpp +// 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) + ")."); + } +} + +// 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 *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); + if (!OclocLibrary) + throw sycl::exception(make_error_code(errc::build), + "Unable to load ocloc library " + OclocLibraryName); + + checkOclocLibrary(OclocLibrary); + + return OclocLibrary; +} + +static std::vector +OpenCLC_to_SPIRV(const std::string &Source, + const std::vector &UserArgs) { + 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) { + void *OclocLibrary = 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; + byte **Outputs = nullptr; + uint64_t *OutputLengths = nullptr; + char **OutputNames = nullptr; + + const byte *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) + std::vector 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 = std::vector(Outputs[I], Outputs[I] + OutputLengths[I]); + } else if (!strcmp(OutputNames[I], "stdout.log")) { + CompileLog = std::string(reinterpret_cast(Outputs[I])); + } + } + + std::cout << "Compilation Log: " << CompileLog << std::endl; + + // 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 \ No newline at end of file diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index a622d9c0c5c7c..d18dc79a9a96b 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -11,7 +11,7 @@ #include #include -// #include +#include #include @@ -397,6 +397,9 @@ exe_kb build(source_kb &SourceKB, const property_list &PropList) { backend BE = SourceKB.get_backend(); + std::vector spirv = detail::OpenCLC_to_SPIRV(s, flags); + std::cout << "spirv byte count: " << spirv.size() << std::endl; + // CP fake code to compile for the nonce. This constructor is empty. std::shared_ptr KBImpl = std::make_shared(SourceKB); From ca4913b1839f76d4d6680757286d97d8c602f67b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 27 Sep 2023 15:47:16 -0700 Subject: [PATCH 05/37] interim. UR sources have been moved. --- .../kernel_compiler/kernel_compiler_opencl.hpp | 18 +++++++++--------- sycl/source/kernel_bundle.cpp | 13 ++++--------- 2 files changed, 13 insertions(+), 18 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index b7ad41b17d19c..bada06fdff430 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -127,18 +127,18 @@ OpenCLC_to_SPIRV(const std::string &Source, // gather the results ( the SpirV and the Log) std::vector 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) { + 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 = std::vector(Outputs[I], Outputs[I] + OutputLengths[I]); - } else if (!strcmp(OutputNames[I], "stdout.log")) { - CompileLog = std::string(reinterpret_cast(Outputs[I])); + SpirV = std::vector(Outputs[i], Outputs[i] + OutputLengths[i]); + } else if (!strcmp(OutputNames[i], "stdout.log")) { + CompileLog = std::string(reinterpret_cast(Outputs[i])); } } - - std::cout << "Compilation Log: " << CompileLog << std::endl; + // std::cout << "Compile Log: " << std::endl << CompileLog << std::endl << + // "=============" << std::endl; // Try to free memory before reporting possible error. decltype(::oclocFreeOutput) *OclocFreeOutputFunc = diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index d18dc79a9a96b..3fadcdaeaa3e8 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -385,19 +385,14 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, exe_kb build(source_kb &SourceKB, const property_list &PropList) { - // CP gross test code. Using existing OnlineCompiler as placeholder. + // CP gross test code. std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); - // sycl::ext::intel::experimental::online_compiler - // Compiler; std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; - std::string s = sourceImpl->Source; - std::cout << "sourcey: " << s << std::endl; - // std::vector SpirVec = Compiler.compile(sourceImpl->Source, flags); - backend BE = SourceKB.get_backend(); - - std::vector spirv = detail::OpenCLC_to_SPIRV(s, flags); + // if successful, the log is empty. if failed, throws an error with the + // compilation log. + std::vector spirv = detail::OpenCLC_to_SPIRV(sourceImpl->Source, flags); std::cout << "spirv byte count: " << spirv.size() << std::endl; // CP fake code to compile for the nonce. This constructor is empty. From 26760b2ed60eb13fb1891f81cf300a72097ebe0d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 28 Sep 2023 15:00:16 -0700 Subject: [PATCH 06/37] attempt to leverage createSpirProgram(). Not working --- sycl/source/CMakeLists.txt | 1 + sycl/source/detail/kernel_bundle_impl.hpp | 40 +++++ .../kernel_compiler_opencl.cpp | 170 ++++++++++++++++++ .../kernel_compiler_opencl.hpp | 155 ++-------------- sycl/source/kernel_bundle.cpp | 16 +- 5 files changed, 234 insertions(+), 148 deletions(-) create mode 100644 sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 1bde471cd67ac..6563d4e6c5739 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/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index f13e6af522e86..6b25897dc6de9 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -10,6 +10,7 @@ #include #include +#include #include #include #include @@ -18,6 +19,8 @@ #include #include + + #include #include #include @@ -332,6 +335,7 @@ class kernel_bundle_impl { // oneapi ext kernel_compiler // construct exe from source kb. + // CP - TEMPORARY EMPTY kernel_bundle_impl( const kernel_bundle &SourceBundle) : MContext(SourceBundle.get_context()), @@ -340,6 +344,42 @@ class kernel_bundle_impl { // sourceImpl = getSyclObjImpl(SourceBundle); } + //std::shared_ptr + void lets_do_this(){ + assert(MState == bundle_state::ext_oneapi_source); + + // CP temp + std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; + + // if successful, the log is empty. if failed, throws an error with the + // compilation log. + //std::vector + auto spirv = syclex::detail::OpenCLC_to_SPIRV(this->Source, flags); + std::cout << "spirv byte count: " << spirv.size() << std::endl; + + // copy/paste from program_manager.cpp::createSpirvProgram() + using ContextImplPtr = std::shared_ptr; + sycl::detail::pi::PiProgram Program = nullptr; + ContextImplPtr ContextImpl = getSyclObjImpl(MContext); + const PluginPtr &Plugin = ContextImpl->getPlugin(); + Plugin->call(ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &Program); + + using exe_kb = kernel_bundle; + sycl::backend Backend = get_backend(); + + std::shared_ptr ExeImpl = sycl::detail::make_kernel_bundle(detail::pi::cast(Program), MContext, bundle_state::executable, Backend); + std::vector kIDs = ExeImpl->get_kernel_ids(); + std::cout << "kernel_ids size: " << kIDs.size() << std::endl; // 0 + + //exe_kb ExecKB = make_kernel_bundle( { Program, sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + + // can't do this here. Could maybe do it in level_zero PI . + // sycl::kernel_bundle SyclKB = + // sycl::make_kernel_bundle( + // {Program, sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + } + bool empty() const noexcept { return MDeviceImages.empty(); } backend get_backend() const noexcept { 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..5949a31c9d372 --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -0,0 +1,170 @@ +//==-- 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 { + +// copy/pasta from online_compiler.cpp +// 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) + ")."); + } +} + +// 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 *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); + if (!OclocLibrary) + throw sycl::exception(make_error_code(errc::build), + "Unable to load ocloc library " + OclocLibraryName); + + checkOclocLibrary(OclocLibrary); + + return OclocLibrary; +} + +spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &UserArgs) { + 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) { + void *OclocLibrary = 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; + //std::byte **Outputs = nullptr; + uint8_t **Outputs = nullptr; + uint64_t *OutputLengths = nullptr; + char **OutputNames = nullptr; + + //const std::byte *Sources[] = {reinterpret_cast(Source.c_str())}; + 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")) { + CompileLog = std::string(reinterpret_cast(Outputs[i])); + } + } + // std::cout << "Compile Log: " << std::endl << CompileLog << std::endl << + // "=============" << std::endl; + + // 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 \ No newline at end of file diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index bada06fdff430..d80a7e1055f68 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -6,161 +6,28 @@ // //===----------------------------------------------------------------------===// + + #pragma once -#include "../online_compiler/ocloc_api.h" +#include +#include // for __SYCL_EXPORT +#include -#include // for std::accumulate +#include +#include namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { -// copy/pasta from online_compiler.cpp -// 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) + ")."); - } -} - -// 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 *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); - if (!OclocLibrary) - throw sycl::exception(make_error_code(errc::build), - "Unable to load ocloc library " + OclocLibraryName); - - checkOclocLibrary(OclocLibrary); - - return OclocLibrary; -} - -static std::vector -OpenCLC_to_SPIRV(const std::string &Source, - const std::vector &UserArgs) { - 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) { - void *OclocLibrary = 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; - byte **Outputs = nullptr; - uint64_t *OutputLengths = nullptr; - char **OutputNames = nullptr; - - const byte *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) - std::vector 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 = std::vector(Outputs[i], Outputs[i] + OutputLengths[i]); - } else if (!strcmp(OutputNames[i], "stdout.log")) { - CompileLog = std::string(reinterpret_cast(Outputs[i])); - } - } - // std::cout << "Compile Log: " << std::endl << CompileLog << std::endl << - // "=============" << std::endl; - - // 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; -} +//using spirv_vec_t = std::vector; +using spirv_vec_t = std::vector; +spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &UserArgs); } // namespace detail } // namespace ext::oneapi::experimental + } // namespace _V1 } // namespace sycl \ No newline at end of file diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 3fadcdaeaa3e8..d67e623ac90d3 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -11,8 +11,6 @@ #include #include -#include - #include namespace sycl { @@ -387,6 +385,7 @@ exe_kb build(source_kb &SourceKB, const property_list &PropList) { // CP gross test code. std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); +/* std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; @@ -395,9 +394,18 @@ exe_kb build(source_kb &SourceKB, const property_list &PropList) { std::vector spirv = detail::OpenCLC_to_SPIRV(sourceImpl->Source, flags); std::cout << "spirv byte count: " << spirv.size() << std::endl; + // copy/paste from program_manager.cpp::createSpirvProgram() + using ContextImplPtr = std::shared_ptr; + sycl::detail::pi::PiProgram Program = nullptr; + ContextImplPtr Context = getSyclObjImpl(SourceKB.get_context()); + const PluginPtr &Plugin = Context->getPlugin(); + Plugin->call(Context->getHandleRef(), spirv.data(), spirv.size(), &Program); + */ + // CP fake code to compile for the nonce. This constructor is empty. - std::shared_ptr KBImpl = - std::make_shared(SourceKB); + sourceImpl->lets_do_this(); + std::shared_ptr KBImpl = std::make_shared(SourceKB); + //std::shared_ptr KBImpl = sourceImpl->lets_do_this(); return sycl::detail::createSyclObjFromImpl(KBImpl); } From dc9e934ae916ce28f6d99053dfe950c175dcc4af Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 28 Sep 2023 18:42:02 -0700 Subject: [PATCH 07/37] interim checkpoint, needs cleaning --- sycl/source/detail/kernel_bundle_impl.hpp | 95 ++++++++++++++++++----- sycl/source/kernel_bundle.cpp | 7 +- 2 files changed, 79 insertions(+), 23 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 6b25897dc6de9..9df10b7f53e1a 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -9,18 +9,17 @@ #pragma once #include -#include #include +#include #include #include #include #include +#include #include #include #include - - #include #include #include @@ -344,8 +343,7 @@ class kernel_bundle_impl { // sourceImpl = getSyclObjImpl(SourceBundle); } - //std::shared_ptr - void lets_do_this(){ + std::shared_ptr lets_do_this() { assert(MState == bundle_state::ext_oneapi_source); // CP temp @@ -353,31 +351,88 @@ class kernel_bundle_impl { // if successful, the log is empty. if failed, throws an error with the // compilation log. - //std::vector auto spirv = syclex::detail::OpenCLC_to_SPIRV(this->Source, flags); std::cout << "spirv byte count: " << spirv.size() << std::endl; - // copy/paste from program_manager.cpp::createSpirvProgram() + // using ContextImplPtr = std::shared_ptr; + // ContextImplPtr ContextImpl = getSyclObjImpl(MContext); + // const PluginPtr &Plugin = ContextImpl->getPlugin(); + // Plugin->call(ContextImpl->getHandleRef(), + // spirv.data(), spirv.size()); + + // // copy/paste from program_manager.cpp::createSpirvProgram() using ContextImplPtr = std::shared_ptr; - sycl::detail::pi::PiProgram Program = nullptr; + sycl::detail::pi::PiProgram PiProgram = nullptr; ContextImplPtr ContextImpl = getSyclObjImpl(MContext); const PluginPtr &Plugin = ContextImpl->getPlugin(); - Plugin->call(ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &Program); + Plugin->call( + ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &PiProgram); - using exe_kb = kernel_bundle; - sycl::backend Backend = get_backend(); + if (ContextImpl->getBackend() == backend::opencl) + Plugin->call(PiProgram); - std::shared_ptr ExeImpl = sycl::detail::make_kernel_bundle(detail::pi::cast(Program), MContext, bundle_state::executable, Backend); - std::vector kIDs = ExeImpl->get_kernel_ids(); - std::cout << "kernel_ids size: " << kIDs.size() << std::endl; // 0 + for (const auto &SyclDev : MDevices) { + std::cout << "device" << std::endl; + pi::PiDevice Dev = getSyclObjImpl(SyclDev)->getHandleRef(); + Plugin->call( + PiProgram, 1, &Dev, nullptr, nullptr, nullptr); + } - //exe_kb ExecKB = make_kernel_bundle( { Program, sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + // Get the number of kernels in the program. + size_t NumKernels; + Plugin->call( + PiProgram, PI_PROGRAM_INFO_NUM_KERNELS, sizeof(size_t), &NumKernels, + nullptr); + std::cout << "Num Kernels: " << NumKernels << std::endl; + + // Get the kernel names. + size_t KernelNamesSize; + Plugin->call( + PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, 0, nullptr, &KernelNamesSize); + std::cout << "KernelNamesSize: " << KernelNamesSize << std::endl; + + std::string KernelNames(KernelNamesSize, + ' '); // semi-colon delimited list of kernel names. + Plugin->call( + PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, KernelNames.size(), + &KernelNames[0], nullptr); + std::cout << "KernelNames: " << KernelNames << std::endl; + + // Create each kernel. + auto names = detail::split_string(KernelNames, ';'); + for (auto name : names) { + sycl::detail::pi::PiKernel Kernel = nullptr; + Plugin->call(PiProgram, name.c_str(), &Kernel); + } - // can't do this here. Could maybe do it in level_zero PI . - // sycl::kernel_bundle SyclKB = - // sycl::make_kernel_bundle( - // {Program, sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + // make the device image and the kernel_bundl_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); + + // using exe_kb = kernel_bundle; + // sycl::backend Backend = get_backend(); + + // std::shared_ptr ExeImpl = + // sycl::detail::make_kernel_bundle(detail::pi::cast(Program), + // MContext, bundle_state::executable, Backend); std::vector kIDs + // = ExeImpl->get_kernel_ids(); std::cout << "kernel_ids size: " << + // kIDs.size() << std::endl; // 0 + + // //exe_kb ExecKB = make_kernel_bundle( { Program, + // sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + + // // can't do this here. Could maybe do it in level_zero PI . + // // sycl::kernel_bundle SyclKB = + // // sycl::make_kernel_bundle( + // // {Program, sycl::ext::oneapi::level_zero::ownership::keep}, + // MContext); } bool empty() const noexcept { return MDeviceImages.empty(); } diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index d67e623ac90d3..31012e8a7dd06 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -403,9 +403,10 @@ exe_kb build(source_kb &SourceKB, const property_list &PropList) { */ // CP fake code to compile for the nonce. This constructor is empty. - sourceImpl->lets_do_this(); - std::shared_ptr KBImpl = std::make_shared(SourceKB); - //std::shared_ptr KBImpl = sourceImpl->lets_do_this(); + std::shared_ptr KBImpl = sourceImpl->lets_do_this(); + // std::shared_ptr KBImpl = + // std::make_shared(SourceKB); + // std::shared_ptr KBImpl = sourceImpl->lets_do_this(); return sycl::detail::createSyclObjFromImpl(KBImpl); } From c536da2f4f547a64669ed83404ed1b4cff63f026 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 29 Sep 2023 16:20:22 -0700 Subject: [PATCH 08/37] except for the fact that the kernels are crashing, this is great! --- sycl/include/sycl/kernel_bundle.hpp | 28 +++++ sycl/source/detail/common.cpp | 18 ++-- sycl/source/detail/kernel_bundle_impl.hpp | 125 ++++++++++++---------- sycl/source/kernel_bundle.cpp | 32 ++---- 4 files changed, 115 insertions(+), 88 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 84257b76eba72..1201aeb2d1352 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -41,6 +41,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 +177,11 @@ class __SYCL_EXPORT kernel_bundle_plain { bool native_specialization_constant() const noexcept; + bool ext_oneapi_has_kernel(const std::string &name); + + std::shared_ptr + ext_oneapi_get_kernel(const std::string &name); + protected: // \returns a kernel object which represents the kernel identified by // kernel_id passed @@ -341,6 +347,28 @@ class kernel_bundle : public detail::kernel_bundle_plain, 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) { + std::shared_ptr kernelImplPtr = + detail::kernel_bundle_plain::ext_oneapi_get_kernel(name); + return sycl::detail::createSyclObjFromImpl(kernelImplPtr); + } + private: kernel_bundle(detail::KernelBundleImplPtr Impl) : kernel_bundle_plain(std::move(Impl)) {} 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 9df10b7f53e1a..0306a6db94c3d 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -46,6 +46,8 @@ static bool checkAllDevicesHaveAspect(const std::vector &Devices, 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. @@ -325,22 +327,21 @@ class kernel_bundle_impl { } } - // oneapi ext kernel_compiler + // 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 - // construct exe from source kb. - // CP - TEMPORARY EMPTY - kernel_bundle_impl( - const kernel_bundle &SourceBundle) - : MContext(SourceBundle.get_context()), - MDevices(SourceBundle.get_devices()), MState(bundle_state::executable) { - - // sourceImpl = getSyclObjImpl(SourceBundle); + // 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 lets_do_this() { @@ -354,13 +355,7 @@ class kernel_bundle_impl { auto spirv = syclex::detail::OpenCLC_to_SPIRV(this->Source, flags); std::cout << "spirv byte count: " << spirv.size() << std::endl; - // using ContextImplPtr = std::shared_ptr; - // ContextImplPtr ContextImpl = getSyclObjImpl(MContext); - // const PluginPtr &Plugin = ContextImpl->getPlugin(); - // Plugin->call(ContextImpl->getHandleRef(), - // spirv.data(), spirv.size()); - - // // copy/paste from program_manager.cpp::createSpirvProgram() + // see also program_manager.cpp::createSpirvProgram() using ContextImplPtr = std::shared_ptr; sycl::detail::pi::PiProgram PiProgram = nullptr; ContextImplPtr ContextImpl = getSyclObjImpl(MContext); @@ -383,56 +378,82 @@ class kernel_bundle_impl { Plugin->call( PiProgram, PI_PROGRAM_INFO_NUM_KERNELS, sizeof(size_t), &NumKernels, nullptr); + // CP std::cout << "Num Kernels: " << NumKernels << std::endl; // Get the kernel names. size_t KernelNamesSize; Plugin->call( PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, 0, nullptr, &KernelNamesSize); - std::cout << "KernelNamesSize: " << KernelNamesSize << std::endl; - std::string KernelNames(KernelNamesSize, - ' '); // semi-colon delimited list of kernel names. + // semi-colon delimited list of kernel names. + std::string KernelNamesStr(KernelNamesSize, ' '); Plugin->call( - PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, KernelNames.size(), - &KernelNames[0], nullptr); - std::cout << "KernelNames: " << KernelNames << std::endl; - + PiProgram, PI_PROGRAM_INFO_KERNEL_NAMES, KernelNamesStr.size(), + &KernelNamesStr[0], nullptr); + std::vector KernelNames = + detail::split_string(KernelNamesStr, ';'); + // CP + std::cout << "KernelNamesStr: " << KernelNamesStr << std::endl; + + // CP // Create each kernel. - auto names = detail::split_string(KernelNames, ';'); - for (auto name : names) { - sycl::detail::pi::PiKernel Kernel = nullptr; - Plugin->call(PiProgram, name.c_str(), &Kernel); - } + // for (auto Name : KernelNames) { + // sycl::detail::pi::PiKernel Kernel = nullptr; + // Plugin->call(PiProgram, Name.c_str(), + // &Kernel); + // } - // make the device image and the kernel_bundl_impl + // 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(); + } - return std::make_shared(MContext, MDevices, DevImg); + std::shared_ptr 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."); - // using exe_kb = kernel_bundle; - // sycl::backend Backend = get_backend(); + if (!ext_oneapi_has_kernel(Name)) + throw sycl::exception(make_error_code(errc::invalid), + "kernel '" + Name + "' not found in kernel_bundle"); - // std::shared_ptr ExeImpl = - // sycl::detail::make_kernel_bundle(detail::pi::cast(Program), - // MContext, bundle_state::executable, Backend); std::vector kIDs - // = ExeImpl->get_kernel_ids(); std::cout << "kernel_ids size: " << - // kIDs.size() << std::endl; // 0 + 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); + + // CP -- alt candidate + // return make_kernel( + // const context &TargetContext, + // const kernel_bundle &KernelBundle, + // pi_native_handle NativeKernelHandle, bool KeepOwnership, backend + // Backend); - // //exe_kb ExecKB = make_kernel_bundle( { Program, - // sycl::ext::oneapi::level_zero::ownership::keep}, MContext); + const KernelArgMask *ArgMask = nullptr; + std::shared_ptr KernelImpl = std::make_shared( + PiKernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self, + ArgMask); - // // can't do this here. Could maybe do it in level_zero PI . - // // sycl::kernel_bundle SyclKB = - // // sycl::make_kernel_bundle( - // // {Program, sycl::ext::oneapi::level_zero::ownership::keep}, - // MContext); + return KernelImpl; } bool empty() const noexcept { return MDeviceImages.empty(); } @@ -633,15 +654,11 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; - // ext_oneapi_kernel_compiler : Source and Languauge + // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames const syclex::source_language Language = syclex::source_language::opencl; const std::string Source; - - // friend declaration for build(source_kb) is a wee ungainly - friend kernel_bundle - sycl::ext::oneapi::experimental::build( - kernel_bundle &SourceKB, - const property_list &PropList); + std::vector + KernelNames; // only kernel_bundles created from source have this. }; } // namespace detail diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 31012e8a7dd06..fd4855147dcab 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -111,6 +111,15 @@ 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); +} + +std::shared_ptr +kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) { + return impl->ext_oneapi_get_kernel(name, impl); +} + ////////////////////////////////// ///// sycl::detail free functions ////////////////////////////////// @@ -382,31 +391,8 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, ///////////////////////// exe_kb build(source_kb &SourceKB, const property_list &PropList) { - - // CP gross test code. std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); -/* - std::vector flags{"-cl-fast-relaxed-math", - "-cl-finite-math-only"}; - - // if successful, the log is empty. if failed, throws an error with the - // compilation log. - std::vector spirv = detail::OpenCLC_to_SPIRV(sourceImpl->Source, flags); - std::cout << "spirv byte count: " << spirv.size() << std::endl; - - // copy/paste from program_manager.cpp::createSpirvProgram() - using ContextImplPtr = std::shared_ptr; - sycl::detail::pi::PiProgram Program = nullptr; - ContextImplPtr Context = getSyclObjImpl(SourceKB.get_context()); - const PluginPtr &Plugin = Context->getPlugin(); - Plugin->call(Context->getHandleRef(), spirv.data(), spirv.size(), &Program); - */ - - // CP fake code to compile for the nonce. This constructor is empty. std::shared_ptr KBImpl = sourceImpl->lets_do_this(); - // std::shared_ptr KBImpl = - // std::make_shared(SourceKB); - // std::shared_ptr KBImpl = sourceImpl->lets_do_this(); return sycl::detail::createSyclObjFromImpl(KBImpl); } From deb2663751a218fd9f7b3088d9db77e0d8973997 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 29 Sep 2023 17:06:50 -0700 Subject: [PATCH 09/37] end to end kernel compilation and execution is working. Need to bring new property list stuff in, error checking, testing, and other to-do. --- sycl/include/sycl/kernel_bundle.hpp | 7 ++----- sycl/source/detail/kernel_bundle_impl.hpp | 22 +++++++++------------- sycl/source/kernel_bundle.cpp | 3 +-- 3 files changed, 12 insertions(+), 20 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 1201aeb2d1352..feb6c256bde0b 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -179,8 +179,7 @@ class __SYCL_EXPORT kernel_bundle_plain { bool ext_oneapi_has_kernel(const std::string &name); - std::shared_ptr - ext_oneapi_get_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 @@ -364,9 +363,7 @@ class kernel_bundle : public detail::kernel_bundle_plain, template > kernel ext_oneapi_get_kernel(const std::string &name) { - std::shared_ptr kernelImplPtr = - detail::kernel_bundle_plain::ext_oneapi_get_kernel(name); - return sycl::detail::createSyclObjFromImpl(kernelImplPtr); + return detail::kernel_bundle_plain::ext_oneapi_get_kernel(name); } private: diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 0306a6db94c3d..abf664151da61 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -419,9 +419,9 @@ class kernel_bundle_impl { return it != KernelNames.end(); } - std::shared_ptr ext_oneapi_get_kernel( - const std::string &Name, - const std::shared_ptr &Self) { + 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 " @@ -441,19 +441,15 @@ class kernel_bundle_impl { sycl::detail::pi::PiKernel PiKernel = nullptr; Plugin->call(PiProgram, Name.c_str(), &PiKernel); - // CP -- alt candidate - // return make_kernel( - // const context &TargetContext, - // const kernel_bundle &KernelBundle, - // pi_native_handle NativeKernelHandle, bool KeepOwnership, backend - // Backend); + // CP ?? - not sure about this. Investigate + // if (Backend == backend::opencl) + // Plugin->call(PiKernel); - const KernelArgMask *ArgMask = nullptr; std::shared_ptr KernelImpl = std::make_shared( - PiKernel, detail::getSyclObjImpl(MContext), DeviceImageImpl, Self, - ArgMask); + PiKernel, detail::getSyclObjImpl(MContext), Self); - return KernelImpl; + return detail::createSyclObjFromImpl(KernelImpl); + ; } bool empty() const noexcept { return MDeviceImages.empty(); } diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index fd4855147dcab..53e949d6343ac 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -115,8 +115,7 @@ bool kernel_bundle_plain::ext_oneapi_has_kernel(const std::string &name) { return impl->ext_oneapi_has_kernel(name); } -std::shared_ptr -kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) { +kernel kernel_bundle_plain::ext_oneapi_get_kernel(const std::string &name) { return impl->ext_oneapi_get_kernel(name, impl); } From 62375948e00c16272351194363f9f2c3b3eee3fb Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 9 Oct 2023 15:55:44 -0700 Subject: [PATCH 10/37] feature test and property list --- .../sycl/ext/oneapi/properties/property.hpp | 3 +- sycl/include/sycl/kernel_bundle.hpp | 44 +++++++++++++++++++ sycl/source/feature_test.hpp.in | 1 + 3 files changed, 47 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index e6220f8a79e3e..0521bd2acf7ae 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -201,8 +201,9 @@ enum PropKind : uint32_t { RegisterAllocMode = 31, GRFSize = 32, GRFSizeAutomatic = 33, + BuildOptions = 34, // PropKindSize must always be the last value. - PropKindSize = 34, + PropKindSize = 35, }; // 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 feb6c256bde0b..3ccb7d2c61a3c 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -20,6 +20,9 @@ #include // for bundle_state #include // for property_list +#include // PropertyT +#include // build_options + #include // for array #include // for size_t, memcpy #include // for function @@ -770,6 +773,35 @@ build(const kernel_bundle &InputBundle, namespace ext::oneapi::experimental { namespace syclex = sycl::ext::oneapi::experimental; + +///////////////////////// +// PropertyT syclex::build_options +///////////////////////// +struct build_options { + std::vector opts; + build_options(const std::string &name) : opts{name} {} + 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 IsCompileTimeProperty : std::true_type {}; + +} // namespace detail + ///////////////////////// // syclex::is_source_kernel_bundle_supported ///////////////////////// @@ -787,10 +819,22 @@ create_kernel_bundle_from_source(const context &SyclContext, ///////////////////////// // syclex::build(source_kb) => exe_kb ///////////////////////// + __SYCL_EXPORT kernel_bundle build(kernel_bundle &SourceKB, const property_list &PropList = {}); +// kernel_bundle +// build(kernel_bundle &SourceKB, const +// property_list &PropList); + +// template> +// __SYCL_EXPORT kernel_bundle +// build(kernel_bundle &SourceKB, +// PropertyListT props = {}) { +// return build(SourceKB, props); +// } + } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 30a76930ca36f..a4bb177df758d 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -43,6 +43,7 @@ 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_MAX_WORK_GROUP_QUERY 1 #define SYCL_EXT_ONEAPI_ND_RANGE_REDUCTIONS 1 #define SYCL_EXT_ONEAPI_DEFAULT_CONTEXT 1 From 0356b9dbaffe13e5c73ae82f6f4178004448f992 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 9 Oct 2023 17:48:32 -0700 Subject: [PATCH 11/37] properties problemties --- .../sycl/ext/oneapi/properties/property.hpp | 3 +- sycl/include/sycl/kernel_bundle.hpp | 63 +++++++++++++++---- sycl/source/kernel_bundle.cpp | 2 +- 3 files changed, 53 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 0521bd2acf7ae..d96df86ffaf8d 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -202,8 +202,9 @@ enum PropKind : uint32_t { GRFSize = 32, GRFSizeAutomatic = 33, BuildOptions = 34, + BuildLog = 35, // PropKindSize must always be the last value. - PropKindSize = 35, + PropKindSize = 36, }; // 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 3ccb7d2c61a3c..d2b735bed6c5e 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -20,8 +20,9 @@ #include // for bundle_state #include // for property_list -#include // PropertyT -#include // build_options +#include // PropertyT +#include // build_options +#include // and log #include // for array #include // for size_t, memcpy @@ -797,11 +798,44 @@ template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildOptions; }; +template <> +struct IsRuntimeProperty : std::true_type {}; + template <> struct IsCompileTimeProperty : std::true_type {}; } // namespace detail +///////////////////////// +// PropertyT syclex::build_log +///////////////////////// +struct build_log { + std::string *log; + build_log(std::string *logArg) : log(logArg) {} +}; +using build_log_key = build_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::true_type {}; + +} // namespace detail + ///////////////////////// // syclex::is_source_kernel_bundle_supported ///////////////////////// @@ -820,20 +854,23 @@ create_kernel_bundle_from_source(const context &SyclContext, // syclex::build(source_kb) => exe_kb ///////////////////////// -__SYCL_EXPORT kernel_bundle -build(kernel_bundle &SourceKB, - const property_list &PropList = {}); -// kernel_bundle -// build(kernel_bundle &SourceKB, const -// property_list &PropList); - -// template> +// OLD // __SYCL_EXPORT kernel_bundle // build(kernel_bundle &SourceKB, -// PropertyListT props = {}) { -// return build(SourceKB, props); -// } +// const property_list &PropList = {}); + +// forward decl +kernel_bundle build_old(kernel_bundle &SourceKB, const property_list &PropList); + +using wtf_kenneth_t = std::tuple; // <-- is this right? + +template> // <-- properties accepts exactly one template arg, and that s.b. a std::tuple?? +__SYCL_EXPORT kernel_bundle +build(kernel_bundle &SourceKB, + PropertyListT props = {}) { + return build_old(SourceKB, /* props */ property_list {}); +} } // namespace ext::oneapi::experimental diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 53e949d6343ac..5a972f036526b 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -389,7 +389,7 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, // syclex::build(source_kb) => exe_kb ///////////////////////// -exe_kb build(source_kb &SourceKB, const property_list &PropList) { +exe_kb build_old(source_kb &SourceKB, const property_list &PropList) { std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); std::shared_ptr KBImpl = sourceImpl->lets_do_this(); return sycl::detail::createSyclObjFromImpl(KBImpl); From f27fcab83e705dc6a298d7277508704fc5a374ee Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 10 Oct 2023 11:58:52 -0700 Subject: [PATCH 12/37] property support --- sycl/include/sycl/kernel_bundle.hpp | 34 +++++++++++-------- sycl/source/detail/kernel_bundle_impl.hpp | 19 ++++++----- .../kernel_compiler_opencl.cpp | 16 +++++---- .../kernel_compiler_opencl.hpp | 4 ++- sycl/source/kernel_bundle.cpp | 11 ++++-- 5 files changed, 52 insertions(+), 32 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index d2b735bed6c5e..d4c4707f90a8b 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -802,7 +802,7 @@ template <> struct IsRuntimeProperty : std::true_type {}; template <> -struct IsCompileTimeProperty : std::true_type {}; +struct IsCompileTimeProperty : std::false_type {}; } // namespace detail @@ -832,7 +832,7 @@ template <> struct IsRuntimeProperty : std::true_type {}; template <> -struct IsCompileTimeProperty : std::true_type {}; +struct IsCompileTimeProperty : std::false_type {}; } // namespace detail @@ -853,23 +853,29 @@ create_kernel_bundle_from_source(const context &SyclContext, ///////////////////////// // syclex::build(source_kb) => exe_kb ///////////////////////// - - -// OLD -// __SYCL_EXPORT kernel_bundle -// build(kernel_bundle &SourceKB, -// const property_list &PropList = {}); - +namespace detail { // forward decl -kernel_bundle build_old(kernel_bundle &SourceKB, const property_list &PropList); - -using wtf_kenneth_t = std::tuple; // <-- is this right? +__SYCL_EXPORT kernel_bundle +build_from_source(kernel_bundle &SourceKB, + const std::vector &BuildOptions, + std::string *LogPtr); +} // namespace detail -template> // <-- properties accepts exactly one template arg, and that s.b. a std::tuple?? +template >> __SYCL_EXPORT kernel_bundle build(kernel_bundle &SourceKB, PropertyListT props = {}) { - return build_old(SourceKB, /* props */ property_list {}); + std::vector BuildOptionsVec; + std::string *LogPtr = nullptr; + if (props.template has_property()) { + BuildOptionsVec = props.template get_property().opts; + } + if (props.template has_property()) { + LogPtr = props.template get_property().log; + } + return detail::build_from_source(SourceKB, BuildOptionsVec, LogPtr); } } // namespace ext::oneapi::experimental diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index abf664151da61..f86def6ed2263 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -344,15 +344,18 @@ class kernel_bundle_impl { KernelNames = KNames; } - std::shared_ptr lets_do_this() { - assert(MState == bundle_state::ext_oneapi_source); - - // CP temp - std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; + std::shared_ptr + build_from_source(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, flags); + auto spirv = + syclex::detail::OpenCLC_to_SPIRV(this->Source, BuildOptions, LogPtr); std::cout << "spirv byte count: " << spirv.size() << std::endl; // see also program_manager.cpp::createSpirvProgram() @@ -653,8 +656,8 @@ class kernel_bundle_impl { // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames const syclex::source_language Language = syclex::source_language::opencl; const std::string Source; - std::vector - KernelNames; // only kernel_bundles created from source have this. + // 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 index 5949a31c9d372..8e6e4be3b450d 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -71,7 +71,9 @@ void *loadOclocLibrary() { return OclocLibrary; } -spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &UserArgs) { +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"); @@ -130,7 +132,7 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector= 4 && strstr(OutputNames[i], ".spv") != nullptr && @@ -138,11 +140,13 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector(Outputs[i])); + // CompileLog = std::string(reinterpret_cast(Outputs[i])); + const char *LogText = reinterpret_cast(Outputs[i]); + if (LogText != nullptr && LogText[0] != '\0') { + LogPtr->append(LogText); + } } } - // std::cout << "Compile Log: " << std::endl << CompileLog << std::endl << - // "=============" << std::endl; // Try to free memory before reporting possible error. decltype(::oclocFreeOutput) *OclocFreeOutputFunc = @@ -152,7 +156,7 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector; using spirv_vec_t = std::vector; -spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &UserArgs); +spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, + const std::vector &UserArgs, + std::string *LogPtr); } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 5a972f036526b..1b77b2c6e0ecd 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -386,15 +386,20 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, } ///////////////////////// -// syclex::build(source_kb) => exe_kb +// syclex::detail::build_from_source(source_kb) => exe_kb ///////////////////////// +namespace detail { -exe_kb build_old(source_kb &SourceKB, const property_list &PropList) { +exe_kb build_from_source(source_kb &SourceKB, + const std::vector &BuildOptions, + std::string *LogPtr) { std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); - std::shared_ptr KBImpl = sourceImpl->lets_do_this(); + std::shared_ptr KBImpl = + sourceImpl->build_from_source(BuildOptions, LogPtr); return sycl::detail::createSyclObjFromImpl(KBImpl); } +} // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 From 27005a6b0cb9c8324b0774572ed2ceed9665f2d4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 11 Oct 2023 14:40:10 -0700 Subject: [PATCH 13/37] tests, constraints mid-work checkpoint --- sycl/include/sycl/kernel_bundle.hpp | 21 ++- sycl/source/detail/kernel_bundle_impl.hpp | 27 +--- .../kernel_compiler_opencl.cpp | 8 +- sycl/source/kernel_bundle.cpp | 5 +- .../KernelCompiler/kernel_compiler.cpp | 148 ++++++++++++++++++ .../kernel_compiler_constraints.cpp | 31 ++++ 6 files changed, 210 insertions(+), 30 deletions(-) create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler.cpp create mode 100644 sycl/test/extensions/kernel_compiler_constraints.cpp diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index d4c4707f90a8b..0fe64e215f6c9 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -274,6 +274,9 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \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(); } @@ -857,6 +860,7 @@ 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 @@ -866,16 +870,25 @@ template >> __SYCL_EXPORT kernel_bundle build(kernel_bundle &SourceKB, - PropertyListT props = {}) { + const std::vector &Devices, PropertyListT props = {}) { std::vector BuildOptionsVec; std::string *LogPtr = nullptr; - if (props.template has_property()) { + if constexpr (props.template has_property()) { BuildOptionsVec = props.template get_property().opts; } - if (props.template has_property()) { + if constexpr (props.template has_property()) { LogPtr = props.template get_property().log; } - return detail::build_from_source(SourceKB, BuildOptionsVec, LogPtr); + return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); +} + +template >> +__SYCL_EXPORT kernel_bundle +build(kernel_bundle &SourceKB, + PropertyListT props = {}) { + return build(SourceKB, SourceKB.get_devices(), props); } } // namespace ext::oneapi::experimental diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index f86def6ed2263..78c7eea7d3407 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -345,7 +345,8 @@ class kernel_bundle_impl { } std::shared_ptr - build_from_source(const std::vector &BuildOptions, + 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"); @@ -356,7 +357,6 @@ class kernel_bundle_impl { // compilation log. auto spirv = syclex::detail::OpenCLC_to_SPIRV(this->Source, BuildOptions, LogPtr); - std::cout << "spirv byte count: " << spirv.size() << std::endl; // see also program_manager.cpp::createSpirvProgram() using ContextImplPtr = std::shared_ptr; @@ -366,11 +366,9 @@ class kernel_bundle_impl { Plugin->call( ContextImpl->getHandleRef(), spirv.data(), spirv.size(), &PiProgram); - if (ContextImpl->getBackend() == backend::opencl) - Plugin->call(PiProgram); + Plugin->call(PiProgram); - for (const auto &SyclDev : MDevices) { - std::cout << "device" << std::endl; + for (const auto &SyclDev : Devices) { pi::PiDevice Dev = getSyclObjImpl(SyclDev)->getHandleRef(); Plugin->call( PiProgram, 1, &Dev, nullptr, nullptr, nullptr); @@ -381,8 +379,6 @@ class kernel_bundle_impl { Plugin->call( PiProgram, PI_PROGRAM_INFO_NUM_KERNELS, sizeof(size_t), &NumKernels, nullptr); - // CP - std::cout << "Num Kernels: " << NumKernels << std::endl; // Get the kernel names. size_t KernelNamesSize; @@ -396,16 +392,6 @@ class kernel_bundle_impl { &KernelNamesStr[0], nullptr); std::vector KernelNames = detail::split_string(KernelNamesStr, ';'); - // CP - std::cout << "KernelNamesStr: " << KernelNamesStr << std::endl; - - // CP - // Create each kernel. - // for (auto Name : KernelNames) { - // sycl::detail::pi::PiKernel Kernel = nullptr; - // Plugin->call(PiProgram, Name.c_str(), - // &Kernel); - // } // make the device image and the kernel_bundle_impl auto KernelIDs = std::make_shared>(); @@ -444,15 +430,12 @@ class kernel_bundle_impl { sycl::detail::pi::PiKernel PiKernel = nullptr; Plugin->call(PiProgram, Name.c_str(), &PiKernel); - // CP ?? - not sure about this. Investigate - // if (Backend == backend::opencl) - // Plugin->call(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(); } diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index 8e6e4be3b450d..2fff832f745e7 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -132,7 +132,7 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, // gather the results ( the SpirV and the Log) spirv_vec_t SpirV; - // std::string CompileLog; + 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 && @@ -143,7 +143,9 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, // CompileLog = std::string(reinterpret_cast(Outputs[i])); const char *LogText = reinterpret_cast(Outputs[i]); if (LogText != nullptr && LogText[0] != '\0') { - LogPtr->append(LogText); + CompileLog.append(LogText); + if (LogPtr != nullptr) + LogPtr->append(LogText); } } } @@ -156,7 +158,7 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, if (CompileError) throw sycl::exception(build_errc, "ocloc reported compilation errors: {\n" + - *LogPtr + "\n}"); + CompileLog + "\n}"); if (SpirV.empty()) throw sycl::exception(build_errc, diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 1b77b2c6e0ecd..0b4c7c0cb6f9a 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -391,11 +391,14 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, 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(BuildOptions, LogPtr); + sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr); return sycl::detail::createSyclObjFromImpl(KBImpl); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp new file mode 100644 index 0000000000000..e518945284433 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -0,0 +1,148 @@ +//==- 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 +// +//===----------------------------------------------------------------------===// + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#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 + +============= + +*/ + +void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, + int added) { + constexpr int N = 4; + int InputArray[N] = {0, 1, 2, 3}; + 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); + }); + + auto Out = OutputBuf.get_access(); + 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 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(); + exe_kb kbExe2 = syclex::build(kbSrc, devs, + syclex::properties{syclex::build_options{flags}, + syclex::build_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! + } + // any other error will escape and cause the test to fail ( as it should ). +} + +int main() { +#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; +} \ No newline at end of file diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp new file mode 100644 index 0000000000000..dd6a65601531f --- /dev/null +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -0,0 +1,31 @@ +//==- 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 with the new bundle_state::ext_oneapi_source should NOT +// support several member functions. This test checks 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(); + 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(); + +#endif +} \ No newline at end of file From 514d0c1f49431026fdc611a05998b89367ccc0cc Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 11 Oct 2023 16:37:53 -0700 Subject: [PATCH 14/37] constraints and test for such --- sycl/include/sycl/kernel_bundle.hpp | 25 ++++++++++- .../kernel_compiler_constraints.cpp | 45 ++++++++++++++++++- 2 files changed, 67 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 0fe64e215f6c9..b4dc4c138a974 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -224,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 { @@ -243,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); } @@ -250,6 +258,9 @@ 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); } @@ -268,6 +279,9 @@ class kernel_bundle : public detail::kernel_bundle_plain, } /// \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(); } @@ -283,6 +297,9 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \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(); } @@ -343,12 +360,18 @@ 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()); } diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index dd6a65601531f..5b2e8088d6f52 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -8,8 +8,8 @@ // RUN: %clangxx -fsyntax-only -fsycl -Xclang -verify -Xclang -verify-ignore-unexpected=note %s -// kernel_bundles with the new bundle_state::ext_oneapi_source should NOT -// support several member functions. This test checks that +// kernel_bundles sporting the new bundle_state::ext_oneapi_source should NOT +// support several member functions. This test confirms that. #include @@ -21,11 +21,52 @@ int main() { 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@../include/sycl/kernel_bundle.hpp:* {{no matching member function for call to 'has_kernel'}} + kbSrc.has_kernel(); + + // expected-error@../include/sycl/kernel_bundle.hpp:* {{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(); + #endif } \ No newline at end of file From e1cd7d9f1d158bae9aaad82ed1c5834e65e79bbb Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 11 Oct 2023 21:27:38 -0700 Subject: [PATCH 15/37] linux ABI and supported check --- sycl/source/kernel_bundle.cpp | 6 ++++-- sycl/test/abi/sycl_symbols_linux.dump | 5 +++++ 2 files changed, 9 insertions(+), 2 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 0b4c7c0cb6f9a..eeb0a4a0c621d 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -368,8 +368,10 @@ 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) { - // TODO - maybe return false? - return true; + // at the moment, OpenCL is the only language supported + // and it's support is limited to the opencl and level_zero backends. + return (BE == sycl::backend::ext_oneapi_level_zero) || + (BE == sycl::backend::opencl); } ///////////////////////// diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index a8a3db4b4e245..7eba36f8d9c91 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3711,10 +3711,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 @@ -3977,6 +3980,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 From 4e6d98666ea6608353545f98b7a74ec2d4fb645c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 12 Oct 2023 11:45:26 -0700 Subject: [PATCH 16/37] clang-format --- .../kernel_compiler/kernel_compiler_opencl.cpp | 12 +++++------- .../kernel_compiler/kernel_compiler_opencl.hpp | 7 ++----- sycl/test-e2e/KernelCompiler/kernel_compiler.cpp | 2 +- sycl/test/extensions/kernel_compiler_constraints.cpp | 2 +- 4 files changed, 9 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index 2fff832f745e7..b15eb01b2f6d7 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -6,8 +6,8 @@ // //===----------------------------------------------------------------------===// -#include // getOsLibraryFuncAddress -#include // for make_error_code +#include // getOsLibraryFuncAddress +#include // for make_error_code #include "kernel_compiler_opencl.hpp" @@ -109,13 +109,12 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, CombinedUserArgs.c_str()}; uint32_t NumOutputs = 0; - //std::byte **Outputs = nullptr; uint8_t **Outputs = nullptr; uint64_t *OutputLengths = nullptr; char **OutputNames = nullptr; - //const std::byte *Sources[] = {reinterpret_cast(Source.c_str())}; - const uint8_t *Sources[] = {reinterpret_cast(Source.c_str())}; + const uint8_t *Sources[] = { + reinterpret_cast(Source.c_str())}; const char *SourceName = "main.cl"; const uint64_t SourceLengths[] = {Source.length() + 1}; @@ -140,7 +139,6 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, 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")) { - // CompileLog = std::string(reinterpret_cast(Outputs[i])); const char *LogText = reinterpret_cast(Outputs[i]); if (LogText != nullptr && LogText[0] != '\0') { CompileLog.append(LogText); @@ -173,4 +171,4 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, } // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 -} // namespace sycl \ No newline at end of file +} // namespace sycl diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index a71f40872697d..a8e15bff55809 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -6,12 +6,10 @@ // //===----------------------------------------------------------------------===// - - #pragma once #include -#include // for __SYCL_EXPORT +#include // for __SYCL_EXPORT #include #include @@ -22,7 +20,6 @@ inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { -//using spirv_vec_t = std::vector; using spirv_vec_t = std::vector; spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, const std::vector &UserArgs, @@ -32,4 +29,4 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, } // namespace ext::oneapi::experimental } // namespace _V1 -} // namespace sycl \ No newline at end of file +} // namespace sycl diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index e518945284433..d61897a9dfa50 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -145,4 +145,4 @@ int main() { static_assert(false, "Kernel Compiler feature test macro undefined"); #endif return 0; -} \ No newline at end of file +} diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index 5b2e8088d6f52..b1231045a2335 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -69,4 +69,4 @@ int main() { kbSrc.empty(); #endif -} \ No newline at end of file +} From ba85eca2700a70b99f8b8f35de2a3969cec223f4 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 12 Oct 2023 12:32:33 -0700 Subject: [PATCH 17/37] clang-format? --- sycl/include/sycl/kernel_bundle.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index b4dc4c138a974..9339dc9b90573 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -824,8 +824,7 @@ template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildOptions; }; -template <> -struct IsRuntimeProperty : std::true_type {}; +template <> struct IsRuntimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::false_type {}; From e15820b9b72a3d88aca3abcc0e89c2aa7fd70003 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 12 Oct 2023 12:46:00 -0700 Subject: [PATCH 18/37] clang-format humiliation --- sycl/include/sycl/kernel_bundle.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 9339dc9b90573..6447608a5c458 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -824,7 +824,8 @@ template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildOptions; }; -template <> struct IsRuntimeProperty : std::true_type {}; +template <> +struct IsRuntimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::false_type {}; @@ -853,8 +854,7 @@ template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildLog; }; -template <> -struct IsRuntimeProperty : std::true_type {}; +template <> struct IsRuntimeProperty : std::true_type {}; template <> struct IsCompileTimeProperty : std::false_type {}; From 8f4443890a616ac8762ab080ae23dd17f0015413 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 12 Oct 2023 15:00:02 -0700 Subject: [PATCH 19/37] removed unneeded export b.c. win --- sycl/include/sycl/kernel_bundle.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 6447608a5c458..60c28b1c26516 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -890,7 +890,7 @@ build_from_source(kernel_bundle &SourceKB, template >> -__SYCL_EXPORT kernel_bundle +kernel_bundle build(kernel_bundle &SourceKB, const std::vector &Devices, PropertyListT props = {}) { std::vector BuildOptionsVec; @@ -907,7 +907,7 @@ build(kernel_bundle &SourceKB, template >> -__SYCL_EXPORT kernel_bundle +kernel_bundle build(kernel_bundle &SourceKB, PropertyListT props = {}) { return build(SourceKB, SourceKB.get_devices(), props); From 6a4b127e576b27b76f06b764baa7d216cc150ae6 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 12 Oct 2023 15:49:46 -0700 Subject: [PATCH 20/37] deprecation of online_compiler --- .../include/sycl/ext/intel/experimental/online_compiler.hpp | 6 +++++- sycl/test/warnings/sycl_2020_deprecations.cpp | 5 +++++ 2 files changed, 10 insertions(+), 1 deletion(-) 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/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}} From c4bd305c4d2372b0fdb7c0655ac17a020357c391 Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Thu, 12 Oct 2023 16:21:27 -0700 Subject: [PATCH 21/37] windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 1a2a2e42ba08b..6fa750fa7bb48 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 From 45ef189cf667c3ee8535c5f0ccefc9a2ad59c9d0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 13 Oct 2023 10:51:27 -0700 Subject: [PATCH 22/37] overlooked feature test --- sycl/source/feature_test.hpp.in | 1 + sycl/test-e2e/KernelCompiler/kernel_compiler.cpp | 11 +++++++++++ 2 files changed, 12 insertions(+) diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 67a466d522d10..20ab9d93950f2 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -44,6 +44,7 @@ inline namespace _V1 { #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/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index d61897a9dfa50..ba56184367289 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -6,9 +6,16 @@ // //===----------------------------------------------------------------------===// +// REQUIRES: cm-compiler + // 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"===( @@ -138,6 +145,10 @@ void test_error() { } 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(); From 6dd666426229d8098c98191a55c91557beaec2fa Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 13 Oct 2023 11:54:51 -0700 Subject: [PATCH 23/37] discussion lead to small change in '_supported' functionality. --- .../kernel_compiler_opencl.cpp | 35 +++++++++++++++---- .../kernel_compiler_opencl.hpp | 2 ++ sycl/source/kernel_bundle.cpp | 21 ++++++++--- 3 files changed, 48 insertions(+), 10 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index b15eb01b2f6d7..132187406e8f8 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -54,6 +54,8 @@ void checkOclocLibrary(void *OclocLibrary) { } } +static void *OclocLibrary = nullptr; + // load the ocloc shared library, check it. void *loadOclocLibrary() { #ifdef __SYCL_RT_OS_WINDOWS @@ -61,16 +63,36 @@ void *loadOclocLibrary() { #else static const std::string OclocLibraryName = "libocloc.so"; #endif - void *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); - if (!OclocLibrary) - throw sycl::exception(make_error_code(errc::build), - "Unable to load ocloc library " + OclocLibraryName); + 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(OclocLibrary); + 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) { @@ -84,7 +106,8 @@ spirv_vec_t OpenCLC_to_SPIRV(const std::string &Source, // setup Library if (!oclocInvokeHandle) { - void *OclocLibrary = loadOclocLibrary(); + if (OclocLibrary == nullptr) + loadOclocLibrary(); oclocInvokeHandle = sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index a8e15bff55809..457e9003e63e1 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -25,6 +25,8 @@ 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 diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index eeb0a4a0c621d..c1a2128f29c6c 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -8,6 +8,7 @@ #include #include +#include #include #include @@ -370,8 +371,14 @@ using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; 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. - return (BE == sycl::backend::ext_oneapi_level_zero) || - (BE == sycl::backend::opencl); + bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) || + (BE == sycl::backend::opencl); + if ((Language == source_language::opencl) && BE_Acceptable) { + return syclex::detail::OpenCLC_Compilation_Available(); + } + + // otherwise + return false; } ///////////////////////// @@ -380,8 +387,14 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { source_kb create_kernel_bundle_from_source(const context &SyclContext, syclex::source_language Language, const std::string &Source) { - // TODO -- throw errc::invalid if lang is not supported by BE. - // use syclex::is_source_kernel_bundle_supported(BE, Lang) + // 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); From 03ced5ebd8d2ce37db5be727c7e68ea0adf49c62 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Nov 2023 15:43:49 -0700 Subject: [PATCH 24/37] small spec changes --- sycl/include/sycl/kernel_bundle.hpp | 45 +++++++++++-------- .../KernelCompiler/kernel_compiler.cpp | 6 +-- 2 files changed, 29 insertions(+), 22 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 60c28b1c26516..b934311577d25 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -799,14 +799,13 @@ build(const kernel_bundle &InputBundle, } namespace ext::oneapi::experimental { -namespace syclex = sycl::ext::oneapi::experimental; ///////////////////////// // PropertyT syclex::build_options ///////////////////////// struct build_options { std::vector opts; - build_options(const std::string &name) : opts{name} {} + build_options(const std::string &optsArg) : opts{optsArg} {} build_options(const std::vector &optsArg) : opts(optsArg) {} }; using build_options_key = build_options; @@ -820,44 +819,51 @@ struct is_property_key_of struct PropertyToKind { +template <> +struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildOptions; }; template <> -struct IsRuntimeProperty : std::true_type {}; +struct IsRuntimeProperty + : std::true_type {}; template <> -struct IsCompileTimeProperty : std::false_type {}; +struct IsCompileTimeProperty + : std::false_type {}; } // namespace detail ///////////////////////// -// PropertyT syclex::build_log +// PropertyT syclex::save_log ///////////////////////// -struct build_log { +struct save_log { std::string *log; - build_log(std::string *logArg) : log(logArg) {} + save_log(std::string *logArg) : log(logArg) {} }; -using build_log_key = build_log; +using save_log_key = save_log; -template <> struct is_property_key : std::true_type {}; +template <> struct is_property_key : std::true_type {}; template <> -struct is_property_key_of> : std::true_type {}; namespace detail { -template <> struct PropertyToKind { +template <> +struct PropertyToKind { static constexpr PropKind Kind = PropKind::BuildLog; }; -template <> struct IsRuntimeProperty : std::true_type {}; +template <> +struct IsRuntimeProperty + : std::true_type {}; template <> -struct IsCompileTimeProperty : std::false_type {}; +struct IsCompileTimeProperty + : std::false_type {}; } // namespace detail @@ -871,9 +877,10 @@ __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, // syclex::create_kernel_bundle_from_source ///////////////////////// __SYCL_EXPORT kernel_bundle -create_kernel_bundle_from_source(const context &SyclContext, - syclex::source_language Language, - const std::string &Source); +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 @@ -898,8 +905,8 @@ build(kernel_bundle &SourceKB, if constexpr (props.template has_property()) { BuildOptionsVec = props.template get_property().opts; } - if constexpr (props.template has_property()) { - LogPtr = props.template get_property().log; + if constexpr (props.template has_property()) { + LogPtr = props.template get_property().log; } return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index ba56184367289..51b601285594b 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -98,9 +98,9 @@ void test_build_and_run() { std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; std::vector devs = kbSrc.get_devices(); - exe_kb kbExe2 = syclex::build(kbSrc, devs, - syclex::properties{syclex::build_options{flags}, - syclex::build_log{&log}}); + 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"); From 9a7e004fd0359e82a5feff09188d9222a1801418 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 2 Nov 2023 16:48:36 -0700 Subject: [PATCH 25/37] some reviewer feedback --- sycl/include/sycl/kernel_bundle.hpp | 18 ++++++++++++++---- .../kernel_compiler/kernel_compiler_opencl.cpp | 2 +- .../kernel_compiler/kernel_compiler_opencl.hpp | 2 +- .../KernelCompiler/kernel_compiler.cpp | 12 +++++++++--- 4 files changed, 25 insertions(+), 9 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index b934311577d25..76ad8b81dfc28 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -267,13 +267,18 @@ class kernel_bundle : public detail::kernel_bundle_plain, /// \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); } @@ -322,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); } @@ -341,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 = diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index 132187406e8f8..9cc8012b6829d 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -1,4 +1,4 @@ -//==-- kernel_compiler_ opencl.cpp OpenCL kernel compilation support -==// +//==-- 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. diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp index 457e9003e63e1..8826cf5ea27a7 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.hpp @@ -1,4 +1,4 @@ -//==-- kernel_compiler_ opencl.hpp OpenCL kernel compilation support -==// +//==-- 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. diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index 51b601285594b..49991f143b59c 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -66,7 +66,7 @@ void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, CGH.parallel_for(sycl::range<1>{N}, Kernel); }); - auto Out = OutputBuf.get_access(); + sycl::host_accessor Out{OutputBuf}; for (int I = 0; I < N; I++) assert(Out[I] == ((I * multiplier) + added)); } @@ -82,8 +82,8 @@ void test_build_and_run() { 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 SOURCE kernel " - "bundle extension: " + std::cout << "Apparently this backend does not support OpenCL C source " + "kernel bundle extension: " << ctx.get_backend() << std::endl; return; } @@ -98,6 +98,11 @@ void test_build_and_run() { 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}}); @@ -140,6 +145,7 @@ void test_error() { 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 ). } From ddc021ff7aad49fc6b2b435e108ebcf3707ba601 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 3 Nov 2023 11:09:17 -0700 Subject: [PATCH 26/37] interim --- sycl/test/extensions/kernel_compiler_constraints.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index dc472bf953da8..b68f7b07c99a5 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -68,5 +68,13 @@ int main() { // 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"}; + // syclex::usm_kind + syclex::build(kbSrc, syclex::properties{ + syclex::build_options{flags}, syclex::save_log{&log}, + syclex::usm_kind}); + #endif } From 4ca2414a083fe84caf2cd6827afa4d6bb1748be5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 3 Nov 2023 13:49:47 -0700 Subject: [PATCH 27/37] checkpoint - having to switch --- .../sycl/ext/oneapi/properties/properties.hpp | 13 +++++++++++++ sycl/include/sycl/kernel_bundle.hpp | 14 +++++++++++++- .../extensions/kernel_compiler_constraints.cpp | 6 ++++++ 3 files changed, 32 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index 2fcc5f7004a5d..cd09d7c509073 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -241,6 +241,19 @@ struct ValueOrDefault< } }; +// CP - DELETE +// template +// struct all_are_property_keys_of; + +// template +// struct all_are_property_keys_of>> : std::true_type +// {}; + +// template +// struct all_are_property_keys_of>> +// : std::bool_constant::value && +// PropTs::template all_are_property_keys_of::value>::value> {}; + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 76ad8b81dfc28..d2e499b89d211 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -902,11 +902,23 @@ build_from_source(kernel_bundle &SourceKB, const std::vector &Devices, const std::vector &BuildOptions, std::string *LogPtr); + +template constexpr bool AllPropsSupported() { + constexpr bool KeyOf[] = { + is_property_key_of, + PropTs>::value...}; + for (size_t i = 0; i < sizeof...(PropTs); i++) + if (!KeyOf[i]) + return false; + return true; +} + } // namespace detail template >> + typename = std::enable_if_t && + detail::AllPropsSupported>> kernel_bundle build(kernel_bundle &SourceKB, const std::vector &Devices, PropertyListT props = {}) { diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index b68f7b07c99a5..18a09e9a20a0f 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -71,6 +71,12 @@ int main() { std::string log; std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; + + syclex::build(kbSrc); + + syclex::build(kbSrc, + syclex::properties{syclex::usm_kind}); + // syclex::usm_kind syclex::build(kbSrc, syclex::properties{ syclex::build_options{flags}, syclex::save_log{&log}, From 7fb648ea12e6a6d7f07378cc91fb9e4ec0fb7a8e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Nov 2023 09:45:52 -0800 Subject: [PATCH 28/37] interim --- .../sycl/ext/oneapi/properties/properties.hpp | 35 ++++++++++++------- sycl/include/sycl/kernel_bundle.hpp | 24 +++++-------- .../kernel_compiler_constraints.cpp | 7 +++- 3 files changed, 37 insertions(+), 29 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index cd09d7c509073..0c0eff53f5813 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -241,18 +241,29 @@ struct ValueOrDefault< } }; -// CP - DELETE -// template -// struct all_are_property_keys_of; - -// template -// struct all_are_property_keys_of>> : std::true_type -// {}; - -// template -// struct all_are_property_keys_of>> -// : std::bool_constant::value && -// PropTs::template all_are_property_keys_of::value>::value> {}; +template struct all_props_ok; + +template +struct all_props_ok : std::true_type {}; + +template +struct all_props_ok + : std::true_type {}; + +template +struct all_props_ok>> + : std::bool_constant< + ext::oneapi::experimental::is_property_key_of::value> { +}; + +template +struct all_props_ok< + SyclT, ext::oneapi::experimental::properties>> + : std::bool_constant< + ext::oneapi::experimental::is_property_key_of::value && + all_props_ok()> {}; } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index d2e499b89d211..2866a0bc70306 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -903,22 +903,15 @@ build_from_source(kernel_bundle &SourceKB, const std::vector &BuildOptions, std::string *LogPtr); -template constexpr bool AllPropsSupported() { - constexpr bool KeyOf[] = { - is_property_key_of, - PropTs>::value...}; - for (size_t i = 0; i < sizeof...(PropTs); i++) - if (!KeyOf[i]) - return false; - return true; -} - } // namespace detail -template && - detail::AllPropsSupported>> +template < + typename PropertyListT = detail::empty_properties_t, + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_ok, + PropertyListT>::value>> + kernel_bundle build(kernel_bundle &SourceKB, const std::vector &Devices, PropertyListT props = {}) { @@ -933,8 +926,7 @@ build(kernel_bundle &SourceKB, return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); } -template >> kernel_bundle build(kernel_bundle &SourceKB, diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index 18a09e9a20a0f..4c5013c85c2e3 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -74,10 +74,15 @@ int main() { syclex::build(kbSrc); + // expected-error@sycl/kernel_bundle.hpp:* {{no matching function for call to 'build'}} syclex::build(kbSrc, syclex::properties{syclex::usm_kind}); - // syclex::usm_kind + 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_ok'}} + // expected-error@sycl/kernel_bundle.hpp:* {{no matching function for call to 'build'}} syclex::build(kbSrc, syclex::properties{ syclex::build_options{flags}, syclex::save_log{&log}, syclex::usm_kind}); From a6964d512415b18f2633c3042dd39c770919d2b0 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Nov 2023 09:50:50 -0800 Subject: [PATCH 29/37] all_props_keys_of added --- .../sycl/ext/oneapi/properties/properties.hpp | 16 ++++++++-------- sycl/include/sycl/kernel_bundle.hpp | 12 ++++++------ .../extensions/kernel_compiler_constraints.cpp | 2 +- 3 files changed, 15 insertions(+), 15 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/properties.hpp b/sycl/include/sycl/ext/oneapi/properties/properties.hpp index 0c0eff53f5813..d3cca173c630f 100644 --- a/sycl/include/sycl/ext/oneapi/properties/properties.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/properties.hpp @@ -241,29 +241,29 @@ struct ValueOrDefault< } }; -template struct all_props_ok; +template struct all_props_are_keys_of; template -struct all_props_ok : std::true_type {}; +struct all_props_are_keys_of : std::true_type {}; template -struct all_props_ok +struct all_props_are_keys_of< + SyclT, ext::oneapi::experimental::detail::empty_properties_t> : std::true_type {}; template -struct all_props_ok>> +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_ok< +struct all_props_are_keys_of< SyclT, ext::oneapi::experimental::properties>> : std::bool_constant< ext::oneapi::experimental::is_property_key_of::value && - all_props_ok()> {}; + all_props_are_keys_of()> {}; } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 2866a0bc70306..5d6aad64133af 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -905,12 +905,12 @@ build_from_source(kernel_bundle &SourceKB, } // namespace detail -template < - typename PropertyListT = detail::empty_properties_t, - typename = std::enable_if_t< - is_property_list_v && - detail::all_props_ok, - PropertyListT>::value>> +template && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> kernel_bundle build(kernel_bundle &SourceKB, diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index 4c5013c85c2e3..d33e129787d7b 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -81,7 +81,7 @@ int main() { 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_ok'}} + // expected-error@../include/sycl/ext/oneapi/properties/properties.hpp:* {{too many template arguments for class template 'all_props_are_keys_of'}} // expected-error@sycl/kernel_bundle.hpp:* {{no matching function for call to 'build'}} syclex::build(kbSrc, syclex::properties{ syclex::build_options{flags}, syclex::save_log{&log}, From 6faecb39a43f3746de7c001e727210ad28a6f687 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Nov 2023 13:11:56 -0800 Subject: [PATCH 30/37] new test for OCL capabilities --- .../KernelCompiler/opencl_capabilities.cpp | 170 ++++++++++++++++++ 1 file changed, 170 insertions(+) create mode 100644 sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp new file mode 100644 index 0000000000000..c397e6f987ca7 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -0,0 +1,170 @@ +//==- 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: cm-compiler + +// 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. + +#include + +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 *a, local float *b, int n) { + if (get_local_id(0) == 0) { + for (int i = 0; i < n; i++) + b[i] = i; + } + barrier(CLK_LOCAL_MEM_FENCE); + + bool ok = true; + for (int i = 0; i < n; i++) + ok &= (b[i] == i); + + a[get_global_id(0)] = ok; + } +)==="; + +void test_local_accessor() { + using namespace sycl; + + 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, int 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"); + + size_t multiplier = 2, added = 100; // the scalars submitted to the kernel + constexpr size_t N = 32; + 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 +// ----------------------- + +auto constexpr StructSrc = R"===( +struct pair { + int multiplier; + int 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 { + int multiplier; + int 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; + constexpr size_t N = 32; + 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; +} \ No newline at end of file From a585ab403399bc9ffb98f4036b2ff4fa06ab0235 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Nov 2023 15:19:21 -0800 Subject: [PATCH 31/37] sign --- sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index c397e6f987ca7..b6e35fffdd317 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -91,7 +91,7 @@ void test_usm_pointer_and_scalar() { exe_kb kbExe1 = syclex::build(kbSrc); sycl::kernel usm_kernel = kbExe1.ext_oneapi_get_kernel("usm_kernel"); - size_t multiplier = 2, added = 100; // the scalars submitted to the kernel + int multiplier = 2, added = 100; // the scalars submitted to the kernel constexpr size_t N = 32; int *usmPtr = sycl::malloc_shared(N, q); From 4901c8449a9df25fc17b72968ca7ad792007c7b2 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 6 Nov 2023 16:23:32 -0800 Subject: [PATCH 32/37] constraint updated --- sycl/include/sycl/kernel_bundle.hpp | 6 ++++- .../kernel_compiler_constraints.cpp | 24 ++++++++++++++++--- 2 files changed, 26 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 5d6aad64133af..b2893854a79d4 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -927,7 +927,11 @@ build(kernel_bundle &SourceKB, } template >> + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> kernel_bundle build(kernel_bundle &SourceKB, PropertyListT props = {}) { diff --git a/sycl/test/extensions/kernel_compiler_constraints.cpp b/sycl/test/extensions/kernel_compiler_constraints.cpp index d33e129787d7b..20f58009ee968 100644 --- a/sycl/test/extensions/kernel_compiler_constraints.cpp +++ b/sycl/test/extensions/kernel_compiler_constraints.cpp @@ -71,21 +71,39 @@ int main() { std::string log; std::vector flags{"-cl-fast-relaxed-math", "-cl-finite-math-only"}; - + // OK syclex::build(kbSrc); - // expected-error@sycl/kernel_bundle.hpp:* {{no matching function for call to 'build'}} + // 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@sycl/kernel_bundle.hpp:* {{no matching function for call to 'build'}} + // 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 } From d0c460549c1dc0c14dd0a17dcc20c2b29fa9147e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 7 Nov 2023 19:59:33 -0800 Subject: [PATCH 33/37] reviewer feedback --- .../kernel_compiler_opencl.cpp | 1 - sycl/source/kernel_bundle.cpp | 5 ++-- .../KernelCompiler/opencl_capabilities.cpp | 25 +++++++++++-------- 3 files changed, 17 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp index 9cc8012b6829d..bf308229a6cb8 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_opencl.cpp @@ -21,7 +21,6 @@ inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { -// copy/pasta from online_compiler.cpp // ensures the OclocLibrary has the right version, etc. void checkOclocLibrary(void *OclocLibrary) { void *OclocVersionHandle = diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index c1a2128f29c6c..4b8e2551d4517 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -360,7 +360,6 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { ///////////////////////// namespace ext::oneapi::experimental { -namespace syclex = sycl::ext::oneapi::experimental; using source_kb = kernel_bundle; using exe_kb = kernel_bundle; using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; @@ -374,7 +373,7 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { bool BE_Acceptable = (BE == sycl::backend::ext_oneapi_level_zero) || (BE == sycl::backend::opencl); if ((Language == source_language::opencl) && BE_Acceptable) { - return syclex::detail::OpenCLC_Compilation_Available(); + return detail::OpenCLC_Compilation_Available(); } // otherwise @@ -385,7 +384,7 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { // syclex::create_kernel_bundle_from_source ///////////////////////// source_kb create_kernel_bundle_from_source(const context &SyclContext, - syclex::source_language Language, + 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 diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index b6e35fffdd317..55860d3f4b709 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -24,18 +24,18 @@ using exe_kb = sycl::kernel_bundle; // local accessor // ----------------------- auto constexpr LocalAccCLSource = R"===( - kernel void test_la(global int *a, local float *b, int n) { + 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++) - b[i] = i; + slm[i] = i + get_group_id(0); } barrier(CLK_LOCAL_MEM_FENCE); bool ok = true; for (int i = 0; i < n; i++) - ok &= (b[i] == i); + ok &= (slm[i] == i + get_group_id(0)); - a[get_global_id(0)] = ok; + buf[get_global_id(0)] = ok; } )==="; @@ -76,7 +76,7 @@ void test_local_accessor() { // USM pointer and scalars // ----------------------- auto constexpr USMCLSource = R"===( -__kernel void usm_kernel(__global int *usmPtr, int multiplier, int added) { +__kernel void usm_kernel(__global int *usmPtr, int multiplier, float added) { size_t i = get_global_id(0); usmPtr[i] = (i * multiplier) + added; } @@ -91,7 +91,9 @@ void test_usm_pointer_and_scalar() { exe_kb kbExe1 = syclex::build(kbSrc); sycl::kernel usm_kernel = kbExe1.ext_oneapi_get_kernel("usm_kernel"); - int multiplier = 2, added = 100; // the scalars submitted to the kernel + // the scalars submitted to the kernel + int multiplier = 2; + float added = 100.f; constexpr size_t N = 32; int *usmPtr = sycl::malloc_shared(N, q); @@ -112,12 +114,15 @@ void test_usm_pointer_and_scalar() { // ----------------------- // 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; - int added; + float added; }; __kernel void struct_kernel(__global int *usmPtr, struct pair adjuster) { size_t i = get_global_id(0); @@ -127,7 +132,7 @@ __kernel void struct_kernel(__global int *usmPtr, struct pair adjuster) { struct pair { int multiplier; - int added; + float added; }; void test_struct() { @@ -140,7 +145,7 @@ void test_struct() { sycl::kernel struct_kernel = kbExe1.ext_oneapi_get_kernel("struct_kernel"); pair adjuster; - adjuster.multiplier = 2, adjuster.added = 100; + adjuster.multiplier = 2, adjuster.added = 100.f; constexpr size_t N = 32; int *usmPtr = sycl::malloc_shared(N, q); @@ -167,4 +172,4 @@ int main() { static_assert(false, "KernelCompiler OpenCL feature test macro undefined"); #endif return 0; -} \ No newline at end of file +} From 1ebedbfcff8b34080e993cf2c7f93c9ecc157e23 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 8 Nov 2023 12:05:32 -0800 Subject: [PATCH 34/37] test requirement fix --- sycl/test-e2e/KernelCompiler/kernel_compiler.cpp | 2 +- sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index 49991f143b59c..4457be88f0c32 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -// REQUIRES: cm-compiler +// REQUIRES: ocloc // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index 55860d3f4b709..aa2d00b1cd259 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -6,7 +6,7 @@ // //===----------------------------------------------------------------------===// -// REQUIRES: cm-compiler +// REQUIRES: ocloc // RUN: %{build} -o %t.out // RUN: %{run} %t.out From 8de6bdf918f6ac08fb14ca5876dcff770489a04c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 10 Nov 2023 09:12:02 -0800 Subject: [PATCH 35/37] cl types --- sycl/test-e2e/KernelCompiler/kernel_compiler.cpp | 4 ++-- sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp | 10 +++++----- 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index 4457be88f0c32..3a2e3d35b2a63 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -54,8 +54,8 @@ Build failed with error code: -11 void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel, int multiplier, int added) { constexpr int N = 4; - int InputArray[N] = {0, 1, 2, 3}; - int OutputArray[N] = {}; + sycl::cl_int InputArray[N] = {0, 1, 2, 3}; + sycl::cl_int OutputArray[N] = {}; sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index aa2d00b1cd259..16d97b81e7a86 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -92,8 +92,8 @@ void test_usm_pointer_and_scalar() { sycl::kernel usm_kernel = kbExe1.ext_oneapi_get_kernel("usm_kernel"); // the scalars submitted to the kernel - int multiplier = 2; - float added = 100.f; + sycl::cl_int multiplier = 2; + sycl::cl_float added = 100.f; constexpr size_t N = 32; int *usmPtr = sycl::malloc_shared(N, q); @@ -131,8 +131,8 @@ __kernel void struct_kernel(__global int *usmPtr, struct pair adjuster) { )==="; struct pair { - int multiplier; - float added; + sycl::cl_int multiplier; + sycl::cl_float added; }; void test_struct() { @@ -147,7 +147,7 @@ void test_struct() { pair adjuster; adjuster.multiplier = 2, adjuster.added = 100.f; constexpr size_t N = 32; - int *usmPtr = sycl::malloc_shared(N, q); + sycl::cl_int *usmPtr = sycl::malloc_shared(N, q); q.submit([&](sycl::handler &cgh) { cgh.set_arg(0, usmPtr); From 010479a3bf60e9243af23b668d20b18e8c9f98fc Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 13 Nov 2023 10:04:22 -0800 Subject: [PATCH 36/37] reviewer feedback - ensure test is also run by L0 --- .../KernelCompiler/kernel_compiler.cpp | 10 +++++---- .../KernelCompiler/opencl_capabilities.cpp | 22 ++++++++++++------- 2 files changed, 20 insertions(+), 12 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp index 3a2e3d35b2a63..9bfce76e06eb9 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp @@ -51,14 +51,16 @@ 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; - sycl::cl_int InputArray[N] = {0, 1, 2, 3}; - sycl::cl_int OutputArray[N] = {}; + 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)); + 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)); diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp index 16d97b81e7a86..65dcbe6eddb0a 100644 --- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp +++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp @@ -14,8 +14,15 @@ // 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. -#include +// 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; @@ -40,7 +47,6 @@ auto constexpr LocalAccCLSource = R"===( )==="; void test_local_accessor() { - using namespace sycl; sycl::queue q; sycl::context ctx = q.get_context(); @@ -92,10 +98,10 @@ void test_usm_pointer_and_scalar() { sycl::kernel usm_kernel = kbExe1.ext_oneapi_get_kernel("usm_kernel"); // the scalars submitted to the kernel - sycl::cl_int multiplier = 2; - sycl::cl_float added = 100.f; + cl_int multiplier = 2; + cl_float added = 100.f; constexpr size_t N = 32; - int *usmPtr = sycl::malloc_shared(N, q); + cl_int *usmPtr = sycl::malloc_shared(N, q); q.submit([&](sycl::handler &cgh) { cgh.set_arg(0, usmPtr); @@ -131,8 +137,8 @@ __kernel void struct_kernel(__global int *usmPtr, struct pair adjuster) { )==="; struct pair { - sycl::cl_int multiplier; - sycl::cl_float added; + cl_int multiplier; + cl_float added; }; void test_struct() { @@ -147,7 +153,7 @@ void test_struct() { pair adjuster; adjuster.multiplier = 2, adjuster.added = 100.f; constexpr size_t N = 32; - sycl::cl_int *usmPtr = sycl::malloc_shared(N, q); + cl_int *usmPtr = sycl::malloc_shared(N, q); q.submit([&](sycl::handler &cgh) { cgh.set_arg(0, usmPtr); From ee77b72d086bc96ad6a39c4b8e0a825b719a5422 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 13 Nov 2023 12:48:46 -0800 Subject: [PATCH 37/37] moved specs to experimental and updated their status --- .../sycl_ext_oneapi_kernel_compiler.asciidoc | 14 ++++++-------- ...sycl_ext_oneapi_kernel_compiler_opencl.asciidoc | 14 ++++++-------- 2 files changed, 12 insertions(+), 16 deletions(-) rename sycl/doc/extensions/{proposed => experimental}/sycl_ext_oneapi_kernel_compiler.asciidoc (97%) rename sycl/doc/extensions/{proposed => experimental}/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc (96%) 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