From c28551885c91aa07f6b90cf126ea6818e7044f67 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 9 Apr 2024 10:55:12 -0700 Subject: [PATCH 01/27] checkpoint --- sycl/include/sycl/kernel_bundle_enums.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/kernel_bundle_enums.hpp b/sycl/include/sycl/kernel_bundle_enums.hpp index 936b0de3879f9..fd53f8cd3a740 100644 --- a/sycl/include/sycl/kernel_bundle_enums.hpp +++ b/sycl/include/sycl/kernel_bundle_enums.hpp @@ -20,7 +20,7 @@ enum class bundle_state : char { namespace ext::oneapi::experimental { -enum class source_language : int { opencl = 0, spirv = 1 /* sycl, cuda */ }; +enum class source_language : int { opencl = 0, spirv = 1, sycl = 2 /* cuda */ }; // opencl versions struct cl_version { From 6c2c911b4786cf4d72a2471c2bbf9db3065c2b26 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 9 Apr 2024 16:23:40 -0700 Subject: [PATCH 02/27] prepare to be templated --- sycl/include/sycl/kernel_bundle.hpp | 58 ++++++++++++++++++++++------- sycl/source/kernel_bundle.cpp | 17 ++++----- 2 files changed, 53 insertions(+), 22 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 5bba4735561a2..e46380b47e50b 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -850,36 +850,68 @@ struct is_property_key_of +make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::vector &Bytes); + +__SYCL_EXPORT kernel_bundle +make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::string &Source); + +__SYCL_EXPORT kernel_bundle +build_from_source(kernel_bundle &SourceKB, + const std::vector &Devices, + const std::vector &BuildOptions, + std::string *LogPtr); + +} // namespace detail + ///////////////////////// // syclex::create_kernel_bundle_from_source ///////////////////////// - +template && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> __SYCL_EXPORT kernel_bundle create_kernel_bundle_from_source(const context &SyclContext, source_language Language, - const std::string &Source); + const std::string &Source, + PropertyListT props = {}) { + // handle the props, which are templated. + + return detail::make_kernel_bundle_from_source(SyclContext, Language, Source); +} #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) +template && + detail::all_props_are_keys_of< + kernel_bundle, + PropertyListT>::value>> __SYCL_EXPORT kernel_bundle create_kernel_bundle_from_source(const context &SyclContext, source_language Language, - const std::vector &Bytes); + const std::vector &Bytes, + PropertyListT props = {}) { + // handle the props, which are templated. + + return detail::make_kernel_bundle_from_source(SyclContext, Language, Bytes); +} #endif ///////////////////////// // syclex::build(source_kb) => exe_kb ///////////////////////// -namespace detail { -// forward decl -__SYCL_EXPORT kernel_bundle -build_from_source(kernel_bundle &SourceKB, - const std::vector &Devices, - const std::vector &BuildOptions, - std::string *LogPtr); - -} // namespace detail -template && detail::all_props_are_keys_of< diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 32ec35dbee837..0d9c16ae73bb2 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -388,12 +388,13 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { } ///////////////////////// -// syclex::create_kernel_bundle_from_source +// syclex::detail::create_kernel_bundle_from_source ///////////////////////// +namespace detail { -source_kb create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source) { +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::string &Source) { // TODO: if we later support a "reason" why support isn't present // (like a missing shared library etc.) it'd be nice to include it in // the exception message here. @@ -407,10 +408,9 @@ source_kb create_kernel_bundle_from_source(const context &SyclContext, return sycl::detail::createSyclObjFromImpl(KBImpl); } -source_kb -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes) { +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::vector &Bytes) { backend BE = SyclContext.get_backend(); if (!is_source_kernel_bundle_supported(BE, Language)) throw sycl::exception(make_error_code(errc::invalid), @@ -424,7 +424,6 @@ create_kernel_bundle_from_source(const context &SyclContext, ///////////////////////// // syclex::detail::build_from_source(source_kb) => exe_kb ///////////////////////// -namespace detail { exe_kb build_from_source(source_kb &SourceKB, const std::vector &Devices, From 60d775949c866ea74a234dfb25d7ee495a1e5acc Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 9 Apr 2024 17:51:54 -0700 Subject: [PATCH 03/27] scaffolding --- .../sycl/ext/oneapi/properties/property.hpp | 3 +- sycl/include/sycl/kernel_bundle.hpp | 51 +++++++++++++++---- sycl/source/CMakeLists.txt | 1 + sycl/source/detail/kernel_bundle_impl.hpp | 19 +++++-- .../kernel_compiler/kernel_compiler_sycl.cpp | 34 +++++++++++++ .../kernel_compiler/kernel_compiler_sycl.hpp | 35 +++++++++++++ sycl/source/kernel_bundle.cpp | 27 +++++++--- 7 files changed, 147 insertions(+), 23 deletions(-) create mode 100644 sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp create mode 100644 sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 559dd20c5fe09..081cb419ce000 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -207,8 +207,9 @@ enum PropKind : uint32_t { SingleTaskKernel = 66, IndirectlyCallable = 67, CallsIndirectly = 68, + IncludeFiles = 69, // PropKindSize must always be the last value. - PropKindSize = 69, + PropKindSize = 70, }; struct property_key_base_tag {}; diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index e46380b47e50b..70e3794e5d5f6 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -814,6 +814,27 @@ build(const kernel_bundle &InputBundle, namespace ext::oneapi::experimental { +///////////////////////// +// PropertyT syclex::include_files +///////////////////////// +struct include_files + : detail::run_time_property_key { + include_files(); + include_files(const std::string &name, const std::string &content) { + record.emplace_back(std::make_pair(name, content)); + } + void add(const std::string &name, const std::string &content) { + record.emplace_back(std::make_pair(name, content)); + } + std::vector> record; +}; +using include_files_key = include_files; + +template <> +struct is_property_key_of> + : std::true_type {}; + ///////////////////////// // PropertyT syclex::build_options ///////////////////////// @@ -853,14 +874,16 @@ __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, namespace detail { // forward decls __SYCL_EXPORT kernel_bundle -make_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes); +make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::string &Source, + std::vector> IncludePairsVec); __SYCL_EXPORT kernel_bundle -make_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source); +make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::vector &Bytes, + std::vector> IncludePairsVec); __SYCL_EXPORT kernel_bundle build_from_source(kernel_bundle &SourceKB, @@ -884,9 +907,13 @@ create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::string &Source, PropertyListT props = {}) { - // handle the props, which are templated. + std::vector> IncludePairsVec; + if constexpr (props.template has_property()) { + IncludePairsVec = props.template get_property().record; + } - return detail::make_kernel_bundle_from_source(SyclContext, Language, Source); + return detail::make_kernel_bundle_from_source(SyclContext, Language, Source, + IncludePairsVec); } #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) @@ -901,9 +928,13 @@ create_kernel_bundle_from_source(const context &SyclContext, source_language Language, const std::vector &Bytes, PropertyListT props = {}) { - // handle the props, which are templated. + std::vector> IncludePairsVec; + if constexpr (props.template has_property()) { + IncludePairsVec = props.template get_property().record; + } - return detail::make_kernel_bundle_from_source(SyclContext, Language, Bytes); + return detail::make_kernel_bundle_from_source(SyclContext, Language, Bytes, + IncludePairsVec); } #endif diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index d683f32d16892..a2a790a623941 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -217,6 +217,7 @@ set(SYCL_COMMON_SOURCES "detail/jit_compiler.cpp" "detail/jit_device_binaries.cpp" "detail/kernel_compiler/kernel_compiler_opencl.cpp" + "detail/kernel_compiler/kernel_compiler_sycl.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 55586b6d2b5ac..229a0e5ace365 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 @@ -331,10 +332,13 @@ class kernel_bundle_impl { // oneapi_ext_kernel_compiler // construct from source string - kernel_bundle_impl(const context &Context, syclex::source_language Lang, - const std::string &Src) + kernel_bundle_impl( + const context &Context, syclex::source_language Lang, + const std::string &Src, + std::vector> IncludePairsVec) : MContext(Context), MDevices(Context.get_devices()), - MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src) {} + MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src), + IncludePairs(IncludePairsVec) {} // oneapi_ext_kernel_compiler // construct from source bytes @@ -397,6 +401,11 @@ class kernel_bundle_impl { [](std::byte B) { return static_cast(B); }); return Result; } + if (Language == syclex::source_language::sycl) { + const auto &SourceStr = std::get(this->Source); + return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs, + BuildOptions, LogPtr); + } throw sycl::exception( make_error_code(errc::invalid), "OpenCL C and SPIR-V are the only supported languages at this time"); @@ -710,11 +719,13 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; - // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames + // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames, IncludePairs const syclex::source_language Language = syclex::source_language::opencl; const std::variant> Source; // only kernel_bundles created from source have KernelNames member. std::vector KernelNames; + std::vector> + IncludePairs; }; } // namespace detail diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp new file mode 100644 index 0000000000000..48f676cbf91ba --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -0,0 +1,34 @@ +//==-- 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 // make_error_code + +#include "kernel_compiler_sycl.hpp" + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +using spirv_vec_t = std::vector; + +spirv_vec_t +SYCL_to_SPIRV(const std::string &Source, + std::vector> IncludePairs, + const std::vector &UserArgs, std::string *LogPtr) {} + +bool SYCL_Compilation_Available() { + // check for clang++ clang++.exe icpx icpx.exe on PATH + + return true; +} + +} // 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_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp new file mode 100644 index 0000000000000..cf26601e447d7 --- /dev/null +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -0,0 +1,35 @@ +//==-- kernel_compiler_sycl.hpp SYCL kernel compilation support -==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include // for __SYCL_EXPORT +#include + +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +using spirv_vec_t = std::vector; +spirv_vec_t +SYCL_to_SPIRV(const std::string &Source, + std::vector> IncludePairs, + const std::vector &UserArgs, std::string *LogPtr); + +bool SYCL_Compilation_Available(); + +} // namespace detail +} // namespace ext::oneapi::experimental + +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 0d9c16ae73bb2..8c3f63397cfd4 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -9,6 +9,7 @@ #include #include #include +#include #include #include @@ -375,11 +376,12 @@ 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 (BE_Acceptable) { - // At the moment, OpenCL and SPIR-V are the only supported languages. if (Language == source_language::opencl) { return detail::OpenCLC_Compilation_Available(); } else if (Language == source_language::spirv) { return true; + } else if (Language == source_language::sycl) { + return detail::SYCL_Compilation_Available(); } } @@ -392,9 +394,10 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { ///////////////////////// namespace detail { -source_kb make_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source) { +source_kb make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::string &Source, + std::vector> IncludePairs) { // 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. @@ -403,19 +406,27 @@ source_kb make_kernel_bundle_from_source(const context &SyclContext, throw sycl::exception(make_error_code(errc::invalid), "kernel_bundle creation from source not supported"); + // throw if include not supported? awaiting guidance + // if(!IncludePairs.empty() && is_include_supported(Languuage)){ throw invalid + // } + std::shared_ptr KBImpl = - std::make_shared(SyclContext, Language, Source); + std::make_shared(SyclContext, Language, Source, + IncludePairs); return sycl::detail::createSyclObjFromImpl(KBImpl); } -source_kb make_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes) { +source_kb make_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::vector &Bytes, + std::vector> IncludePairs) { 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"); + // throw if !IncludePairs.empty() ? awaiting guidance. + std::shared_ptr KBImpl = std::make_shared(SyclContext, Language, Bytes); return sycl::detail::createSyclObjFromImpl(KBImpl); From 06deabf4654e56537f28761a26556d3f5c9f686d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 9 Apr 2024 19:06:35 -0700 Subject: [PATCH 04/27] kernel_compiler testing, preliminary --- ...ompiler.cpp => kernel_compiler_opencl.cpp} | 2 +- .../KernelCompiler/kernel_compiler_sycl.cpp | 109 ++++++++++++++++++ 2 files changed, 110 insertions(+), 1 deletion(-) rename sycl/test-e2e/KernelCompiler/{kernel_compiler.cpp => kernel_compiler_opencl.cpp} (98%) create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp similarity index 98% rename from sycl/test-e2e/KernelCompiler/kernel_compiler.cpp rename to sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 356a26f7be2e8..9427442fa7efc 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -1,4 +1,4 @@ -//==- kernel_compiler.cpp --- kernel_compiler extension tests -------------==// +//==- kernel_compiler_opencl.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. diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp new file mode 100644 index 0000000000000..2d3bf5b2679c2 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -0,0 +1,109 @@ +//==- kernel_compiler_sycl.cpp --- kernel_compiler extension tests -------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) +// UNSUPPORTED: accelerator + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include + +auto constexpr AddEmH = R"===( + int AddEm(int a, int b){ + return a + b + 5; + } +)==="; + +auto constexpr SYCLSource = R"===( +#include +#include "AddEm.h" + +SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = AddEm(GId.get(0), 37); +} +)==="; + +void test_1(sycl::queue &Queue, sycl::kernel &Kernel) { + constexpr int Range = 10; + int *usmPtr = sycl::malloc_shared(Range, Queue); + int start = 3; + + sycl::nd_range<1> R1{{Range}, {1}}; + + bool Passa = true; + + memset(usmPtr, 0, Range * sizeof(int)); + Queue.submit([&](sycl::handler &Handler) { + Handler.set_arg(0, usmPtr); + // Handler.set_arg(1, start); + // Handler.set_arg(2, Range); + Handler.parallel_for(R1, Kernel); + }); + Queue.wait(); + + for (int i = 0; i < Range; i++) { + std::cout << usmPtr[i] << " "; + // assert(usmPtr[i] = i + 42); + } + std::cout << std::endl; + + sycl::free(usmPtr, Queue); +} + +void test_build_and_run() { + namespace syclex = sycl::ext::oneapi::experimental; + // this dance avoids a bug on L0, ensuring context is of exactly one device + sycl::device d; + sycl::context ctx{d}; + sycl::queue q{ctx, d}; + + bool ok = syclex::is_source_kernel_bundle_supported( + ctx.get_backend(), syclex::source_language::sycl); + if (!ok) { + std::cout << "Apparently this backend does not support SYCL source " + "kernel bundle extension: " + << ctx.get_backend() << std::endl; + return; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::opencl, SYCLSource); + // compilation of empty prop list, no devices + exe_kb kbExe1 = syclex::build(kbSrc); + + sycl::kernel k = kbExe.ext_oneapi_get_kernel( + "__free_function_ff_cp"); // amend __free_function_ to kernel f name. + + // NOTE THIS NOISE + // sycl::kernel_bundle kb = + // syclexp::build(kb_src, + // syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); + // sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar"); + + // 4 + test_1(q, k); +} + +int main() { + // TODO - awaiting guidance + // #ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_SYCL + // static_assert(false, "KernelCompiler SYCL feature test macro undefined"); + // #endif + +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + test_build_and_run(); + // test_error(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} From 4ea783615af691e9cda501ea48044156d8e72e65 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 11:48:55 -0700 Subject: [PATCH 05/27] working e2e. next: flags, errors, removal of is_source_kernel_bundle_supported, minutiae --- .../kernel_compiler/kernel_compiler_sycl.cpp | 192 +++++++++++++++++- .../KernelCompiler/kernel_compiler_sycl.cpp | 23 ++- 2 files changed, 207 insertions(+), 8 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 48f676cbf91ba..0646ca071a8a9 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -8,6 +8,12 @@ #include // make_error_code +#include +#include +#include +#include +#include + #include "kernel_compiler_sycl.hpp" namespace sycl { @@ -15,12 +21,194 @@ inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { +std::string generate_semi_unique_id() { + // Get the current time as a time_t object + std::time_t now = std::time(nullptr); + + // Convert time_t to a string with format YYYYMMDD_HHMMSS + std::tm *local_time = std::localtime(&now); + std::stringstream ss; + ss << std::put_time(local_time, "%Y%m%d_%H%M%S"); + + // amend with random number + std::random_device rd; + int random_number = rd() % 900 + 100; + ss << "_" << std::setfill('0') << std::setw(3) + << random_number; // Pad with leading zeros + + return ss.str(); +} + +std::filesystem::path prepare_ws(const std::string &id) { + // const std::filesystem::path tmp = std::filesystem::current_path(); + const std::filesystem::path tmp = std::filesystem::temp_directory_path(); + + std::filesystem::path new_directory_path = tmp / id; + + // std::cout << "tmp: " << tmp << " id: " << id << std::endl; + + // will throw an exception ( fs::filesystem_error ) + // should we catch that and change it to errc::build ? + std::filesystem::create_directories(new_directory_path); + // CP + std::cout << "Directory created: " << new_directory_path << std::endl; + + return new_directory_path; +} + +void output_preamble(std::ofstream &os, const std::filesystem::path &file_path, + const std::string &id) { + + os << "/*\n clang++ -fsycl -o " << id << ".bin -fsycl-dump-device-code=./ " + << id << ".cpp \n */" << std::endl; +} + +std::filesystem::path output_cpp(const std::filesystem::path &parent_dir, + const std::string &id, + std::string raw_code_string) { + std::filesystem::path file_path = parent_dir / (id + ".cpp"); + std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); + + if (outfile.is_open()) { + output_preamble(outfile, file_path, id); + outfile << raw_code_string << std::endl; + + // temporarily needed until -c works with -fsycl-dump-spirv + outfile << "int main(){ return 0; }" << std::endl; + + outfile.close(); // Close the file when finished + } else { + throw sycl::exception(sycl::errc::build, + "Failed to open .cpp file for write: " + + file_path.string()); + } + return file_path; +} + +void output_include_files( + const std::filesystem::path &dpath, + std::vector> IncludePairs) { + using pairStrings = std::pair; + for (pairStrings p : IncludePairs) { + std::filesystem::path file_path = dpath / p.first; + std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); + if (outfile.is_open()) { + outfile << p.second << std::endl; + + outfile.close(); + } else { + throw sycl::exception(sycl::errc::build, + "Failed to open include file for write: " + + file_path.string()); + } + } +} + +void invoke_compiler(const std::filesystem::path &fpath, + const std::filesystem::path &dpath, const std::string &id, + const std::vector &UserArgs, + std::string *LogPtr) { + + std::filesystem::path file_path(fpath); + std::filesystem::path parent_dir(dpath); + std::filesystem::path target_path = parent_dir / (id + ".bin"); + std::filesystem::path log_path = parent_dir / "compilation_log.txt"; +#ifdef __WIN32 + std::string compiler = "clang++.exe"; +#else + std::string compiler = "clang++"; +#endif + + // TODO: UserArgs!!! + + std::string command = + compiler + " -fsycl -o " + target_path.make_preferred().string() + + " -fsycl-dump-device-code=" + parent_dir.make_preferred().string() + " " + + file_path.make_preferred().string() + " 2> " + + log_path.make_preferred().string(); + + // CP + std::cout << "command: " << command << std::endl; + + int result = std::system(command.c_str()); + + // Read the log file contents into the log variable + std::string CompileLog; + std::ifstream log_stream; + log_stream.open(log_path); + if (log_stream.is_open()) { + std::stringstream log_buffer; + log_buffer << log_stream.rdbuf(); + CompileLog.append(log_buffer.str()); + if (LogPtr != nullptr) + LogPtr->append(log_buffer.str()); + + // CP + std::cout << "compile log: " << CompileLog << std::endl; + } else if (result == 0 && LogPtr != nullptr) { + // if there was a compilation problem, we want to report that (below) + // not a mere "missing log" error. + throw sycl::exception(sycl::errc::build, + "failure retrieving compilation log"); + } + + if (result != 0) { + throw sycl::exception(sycl::errc::build, + "Compile failure: " + std::to_string(result) + " " + + CompileLog); + } +} + +std::filesystem::path find_spv(const std::filesystem::path &parent_dir, + const std::string &id) { + std::regex pattern_regex(id + R"(.*\.spv)"); + + // Iterate through all files in the directory matching the pattern + for (const auto &entry : std::filesystem::directory_iterator(parent_dir)) { + if (entry.is_regular_file() && + std::regex_match(entry.path().filename().string(), pattern_regex)) { + // Found the matching file + // CP + std::cout << "Matching file found: " << entry.path() << std::endl; + return entry.path(); + } + } + // File not found, throw + throw sycl::exception(sycl::errc::build, "SPIRV output matching " + id + + " missing from " + + parent_dir.filename().string()); +} + +using spirv_vec_t = std::vector; + +spirv_vec_t loadSPIRVFromFile(std::filesystem::path file_name) { + std::ifstream spv_stream(file_name, std::ios::binary); + spv_stream.seekg(0, std::ios::end); + size_t sz = spv_stream.tellg(); + spv_stream.seekg(0); + // std::vector spv(sz); + spirv_vec_t spv(sz); + spv_stream.read(reinterpret_cast(spv.data()), sz); + + return spv; +} + using spirv_vec_t = std::vector; spirv_vec_t -SYCL_to_SPIRV(const std::string &Source, +SYCL_to_SPIRV(const std::string &SYCLSource, std::vector> IncludePairs, - const std::vector &UserArgs, std::string *LogPtr) {} + const std::vector &UserArgs, std::string *LogPtr) { + const std::string id = generate_semi_unique_id(); + const std::filesystem::path parent_dir = prepare_ws(id); + std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource); + output_include_files(parent_dir, IncludePairs); + invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); + std::filesystem::path spv_path = find_spv(parent_dir, id); + return loadSPIRVFromFile(spv_path); + + // throw sycl::exception(sycl::errc::build, "hi"); +} bool SYCL_Compilation_Available() { // check for clang++ clang++.exe icpx icpx.exe on PATH diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 2d3bf5b2679c2..c47937db1564f 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -26,7 +26,8 @@ auto constexpr SYCLSource = R"===( SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) void ff_cp(int *ptr) { - sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + //sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); sycl::id<1> GId = Item.get_global_id(); ptr[GId.get(0)] = AddEm(GId.get(0), 37); } @@ -61,11 +62,20 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel) { void test_build_and_run() { namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + using exe_kb = sycl::kernel_bundle; + + // TODO: remove this dance from other tests // this dance avoids a bug on L0, ensuring context is of exactly one device - sycl::device d; - sycl::context ctx{d}; - sycl::queue q{ctx, d}; + // sycl::device d; + // sycl::context ctx{d}; + // sycl::queue q{ctx, d}; + + sycl::queue q; + sycl::context ctx = q.get_context(); + // TODO: replace is_source_kernel_bundle_supported() with + // device::ext_oneapi_can_compile() bool ok = syclex::is_source_kernel_bundle_supported( ctx.get_backend(), syclex::source_language::sycl); if (!ok) { @@ -76,9 +86,10 @@ void test_build_and_run() { } source_kb kbSrc = syclex::create_kernel_bundle_from_source( - ctx, syclex::source_language::opencl, SYCLSource); + ctx, syclex::source_language::sycl, SYCLSource, + syclex::properties{syclex::include_files{"AddEm.h", AddEmH}}); // compilation of empty prop list, no devices - exe_kb kbExe1 = syclex::build(kbSrc); + exe_kb kbExe = syclex::build(kbSrc); sycl::kernel k = kbExe.ext_oneapi_get_kernel( "__free_function_ff_cp"); // amend __free_function_ to kernel f name. From 54be4d000d647b4861528538fcaacda15ec6d5a1 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 13:35:51 -0700 Subject: [PATCH 06/27] housekeeping before continuing. also, clang-format sux --- sycl/source/detail/kernel_bundle_impl.hpp | 11 ++--- .../kernel_compiler/kernel_compiler_sycl.cpp | 49 ++++++++----------- .../kernel_compiler/kernel_compiler_sycl.hpp | 10 ++-- sycl/source/kernel_bundle.cpp | 18 ++++--- .../KernelCompiler/kernel_compiler_sycl.cpp | 1 + 5 files changed, 43 insertions(+), 46 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 229a0e5ace365..f35cf15bc9ce3 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -330,12 +330,12 @@ class kernel_bundle_impl { } } + using include_pairs_t = + std::vector>; // oneapi_ext_kernel_compiler // construct from source string - kernel_bundle_impl( - const context &Context, syclex::source_language Lang, - const std::string &Src, - std::vector> IncludePairsVec) + kernel_bundle_impl(const context &Context, syclex::source_language Lang, + const std::string &Src, include_pairs_t IncludePairsVec) : MContext(Context), MDevices(Context.get_devices()), MState(bundle_state::ext_oneapi_source), Language(Lang), Source(Src), IncludePairs(IncludePairsVec) {} @@ -724,8 +724,7 @@ class kernel_bundle_impl { const std::variant> Source; // only kernel_bundles created from source have KernelNames member. std::vector KernelNames; - std::vector> - IncludePairs; + include_pairs_t IncludePairs; }; } // namespace detail diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 0646ca071a8a9..b43c7aea2b28f 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -40,16 +40,16 @@ std::string generate_semi_unique_id() { } std::filesystem::path prepare_ws(const std::string &id) { - // const std::filesystem::path tmp = std::filesystem::current_path(); const std::filesystem::path tmp = std::filesystem::temp_directory_path(); std::filesystem::path new_directory_path = tmp / id; - // std::cout << "tmp: " << tmp << " id: " << id << std::endl; + try { + std::filesystem::create_directories(new_directory_path); + } catch (std::filesystem::filesystem_error const &e) { + throw sycl::exception(sycl::errc::build, e.what()); + } - // will throw an exception ( fs::filesystem_error ) - // should we catch that and change it to errc::build ? - std::filesystem::create_directories(new_directory_path); // CP std::cout << "Directory created: " << new_directory_path << std::endl; @@ -85,9 +85,8 @@ std::filesystem::path output_cpp(const std::filesystem::path &parent_dir, return file_path; } -void output_include_files( - const std::filesystem::path &dpath, - std::vector> IncludePairs) { +void output_include_files(const std::filesystem::path &dpath, + include_pairs_t IncludePairs) { using pairStrings = std::pair; for (pairStrings p : IncludePairs) { std::filesystem::path file_path = dpath / p.first; @@ -167,7 +166,6 @@ std::filesystem::path find_spv(const std::filesystem::path &parent_dir, for (const auto &entry : std::filesystem::directory_iterator(parent_dir)) { if (entry.is_regular_file() && std::regex_match(entry.path().filename().string(), pattern_regex)) { - // Found the matching file // CP std::cout << "Matching file found: " << entry.path() << std::endl; return entry.path(); @@ -179,35 +177,30 @@ std::filesystem::path find_spv(const std::filesystem::path &parent_dir, parent_dir.filename().string()); } -using spirv_vec_t = std::vector; - -spirv_vec_t loadSPIRVFromFile(std::filesystem::path file_name) { +spirv_vec_t load_spv_from_file(std::filesystem::path file_name) { std::ifstream spv_stream(file_name, std::ios::binary); spv_stream.seekg(0, std::ios::end); size_t sz = spv_stream.tellg(); spv_stream.seekg(0); - // std::vector spv(sz); spirv_vec_t spv(sz); spv_stream.read(reinterpret_cast(spv.data()), sz); return spv; } -using spirv_vec_t = std::vector; - -spirv_vec_t -SYCL_to_SPIRV(const std::string &SYCLSource, - std::vector> IncludePairs, - const std::vector &UserArgs, std::string *LogPtr) { - const std::string id = generate_semi_unique_id(); - const std::filesystem::path parent_dir = prepare_ws(id); - std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource); - output_include_files(parent_dir, IncludePairs); - invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); - std::filesystem::path spv_path = find_spv(parent_dir, id); - return loadSPIRVFromFile(spv_path); - - // throw sycl::exception(sycl::errc::build, "hi"); +spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, + include_pairs_t IncludePairs, + const std::vector &UserArgs, + std::string *LogPtr) { + // clang-format off + const std::string id = generate_semi_unique_id(); + const std::filesystem::path parent_dir = prepare_ws(id); + std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource); + output_include_files(parent_dir, IncludePairs); + invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); + std::filesystem::path spv_path = find_spv(parent_dir, id); + return load_spv_from_file(spv_path); + // clang-format on } bool SYCL_Compilation_Available() { diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index cf26601e447d7..799f493787a24 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -21,10 +21,12 @@ namespace ext::oneapi::experimental { namespace detail { using spirv_vec_t = std::vector; -spirv_vec_t -SYCL_to_SPIRV(const std::string &Source, - std::vector> IncludePairs, - const std::vector &UserArgs, std::string *LogPtr); +using include_pairs_t = std::vector>; + +spirv_vec_t SYCL_to_SPIRV(const std::string &Source, + include_pairs_t IncludePairs, + const std::vector &UserArgs, + std::string *LogPtr); bool SYCL_Compilation_Available(); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 8c3f63397cfd4..cd35f3dc0d084 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -394,10 +394,12 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { ///////////////////////// namespace detail { -source_kb make_kernel_bundle_from_source( - const context &SyclContext, source_language Language, - const std::string &Source, - std::vector> IncludePairs) { +using include_pairs_t = std::vector>; + +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::string &Source, + include_pairs_t IncludePairs) { // 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. @@ -416,10 +418,10 @@ source_kb make_kernel_bundle_from_source( return sycl::detail::createSyclObjFromImpl(KBImpl); } -source_kb make_kernel_bundle_from_source( - const context &SyclContext, source_language Language, - const std::vector &Bytes, - std::vector> IncludePairs) { +source_kb make_kernel_bundle_from_source(const context &SyclContext, + source_language Language, + const std::vector &Bytes, + include_pairs_t IncludePairs) { backend BE = SyclContext.get_backend(); if (!is_source_kernel_bundle_supported(BE, Language)) throw sycl::exception(make_error_code(errc::invalid), diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index c47937db1564f..fe12fdea59611 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -85,6 +85,7 @@ void test_build_and_run() { return; } + // TODO: replace with device.ext_support_blah_nha source_kb kbSrc = syclex::create_kernel_bundle_from_source( ctx, syclex::source_language::sycl, SYCLSource, syclex::properties{syclex::include_files{"AddEm.h", AddEmH}}); From 1a9cf00f730511835f192956fb4c56be8b4b150b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 16:19:30 -0700 Subject: [PATCH 07/27] flags and testing of such --- .../kernel_compiler/kernel_compiler_sycl.cpp | 37 +++++---- .../KernelCompiler/kernel_compiler_sycl.cpp | 77 ++++++++++++++++--- 2 files changed, 86 insertions(+), 28 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index b43c7aea2b28f..734d6aefd3200 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -50,27 +50,34 @@ std::filesystem::path prepare_ws(const std::string &id) { throw sycl::exception(sycl::errc::build, e.what()); } - // CP - std::cout << "Directory created: " << new_directory_path << std::endl; - return new_directory_path; } +std::string user_args_as_string(const std::vector &UserArgs) { + return std::accumulate(UserArgs.begin(), UserArgs.end(), std::string(""), + [](const std::string &a, const std::string &b) { + return a.empty() ? b : a + " " + b; + }); +} + void output_preamble(std::ofstream &os, const std::filesystem::path &file_path, - const std::string &id) { + const std::string &id, + const std::vector &UserArgs) { - os << "/*\n clang++ -fsycl -o " << id << ".bin -fsycl-dump-device-code=./ " - << id << ".cpp \n */" << std::endl; + os << "/*\n clang++ -fsycl -o " << id << ".bin " + << user_args_as_string(UserArgs) << " -fsycl-dump-device-code=./ " << id + << ".cpp \n */" << std::endl; } std::filesystem::path output_cpp(const std::filesystem::path &parent_dir, const std::string &id, - std::string raw_code_string) { + std::string raw_code_string, + const std::vector &UserArgs) { std::filesystem::path file_path = parent_dir / (id + ".cpp"); std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); if (outfile.is_open()) { - output_preamble(outfile, file_path, id); + output_preamble(outfile, file_path, id, UserArgs); outfile << raw_code_string << std::endl; // temporarily needed until -c works with -fsycl-dump-spirv @@ -118,17 +125,13 @@ void invoke_compiler(const std::filesystem::path &fpath, std::string compiler = "clang++"; #endif - // TODO: UserArgs!!! - std::string command = - compiler + " -fsycl -o " + target_path.make_preferred().string() + + compiler + " -fsycl -o " + target_path.make_preferred().string() + " " + + user_args_as_string(UserArgs) + " -fsycl-dump-device-code=" + parent_dir.make_preferred().string() + " " + file_path.make_preferred().string() + " 2> " + log_path.make_preferred().string(); - // CP - std::cout << "command: " << command << std::endl; - int result = std::system(command.c_str()); // Read the log file contents into the log variable @@ -142,8 +145,6 @@ void invoke_compiler(const std::filesystem::path &fpath, if (LogPtr != nullptr) LogPtr->append(log_buffer.str()); - // CP - std::cout << "compile log: " << CompileLog << std::endl; } else if (result == 0 && LogPtr != nullptr) { // if there was a compilation problem, we want to report that (below) // not a mere "missing log" error. @@ -166,8 +167,6 @@ std::filesystem::path find_spv(const std::filesystem::path &parent_dir, for (const auto &entry : std::filesystem::directory_iterator(parent_dir)) { if (entry.is_regular_file() && std::regex_match(entry.path().filename().string(), pattern_regex)) { - // CP - std::cout << "Matching file found: " << entry.path() << std::endl; return entry.path(); } } @@ -195,7 +194,7 @@ spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, // clang-format off const std::string id = generate_semi_unique_id(); const std::filesystem::path parent_dir = prepare_ws(id); - std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource); + std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource, UserArgs); output_include_files(parent_dir, IncludePairs); invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); std::filesystem::path spv_path = find_spv(parent_dir, id); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index fe12fdea59611..950ea38d24d6a 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -26,13 +26,28 @@ auto constexpr SYCLSource = R"===( SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) void ff_cp(int *ptr) { - //sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); - sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>() + sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); + sycl::id<1> GId = Item.get_global_id(); ptr[GId.get(0)] = AddEm(GId.get(0), 37); } )==="; +auto constexpr BadSource = R"===( +#include + +SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_cp(int *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id() + no semi colon !! + ptr[GId.get(0)] = GId.get(0) + 41; +} +)==="; + void test_1(sycl::queue &Queue, sycl::kernel &Kernel) { constexpr int Range = 10; int *usmPtr = sycl::malloc_shared(Range, Queue); @@ -76,8 +91,10 @@ void test_build_and_run() { // TODO: replace is_source_kernel_bundle_supported() with // device::ext_oneapi_can_compile() - bool ok = syclex::is_source_kernel_bundle_supported( - ctx.get_backend(), syclex::source_language::sycl); + // bool ok = syclex::is_source_kernel_bundle_supported(ctx.get_backend(), + // syclex::source_language::sycl); + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); if (!ok) { std::cout << "Apparently this backend does not support SYCL source " "kernel bundle extension: " @@ -85,15 +102,32 @@ void test_build_and_run() { return; } - // TODO: replace with device.ext_support_blah_nha + // create from source source_kb kbSrc = syclex::create_kernel_bundle_from_source( ctx, syclex::source_language::sycl, SYCLSource, syclex::properties{syclex::include_files{"AddEm.h", AddEmH}}); + + // double check kernel_bundle.get_source() / get_backend() + sycl::context ctxRes = kbSrc.get_context(); + assert(ctxRes == ctx); + sycl::backend beRes = kbSrc.get_backend(); + assert(beRes == ctx.get_backend()); + // compilation of empty prop list, no devices - exe_kb kbExe = syclex::build(kbSrc); + exe_kb kbExe1 = syclex::build(kbSrc); + + // compilation with props and devices + std::string log; + std::vector flags{"-g", "-fno-fast-math"}; + std::vector devs = kbSrc.get_devices(); + exe_kb kbExe2 = syclex::build( + kbSrc, devs, + syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); + assert(log.find("warning: 'this_nd_item<1>' is deprecated") != + std::string::npos); - sycl::kernel k = kbExe.ext_oneapi_get_kernel( - "__free_function_ff_cp"); // amend __free_function_ to kernel f name. + // amend __free_function_ to kernel f name. + sycl::kernel k = kbExe2.ext_oneapi_get_kernel("__free_function_ff_cp"); // NOTE THIS NOISE // sycl::kernel_bundle kb = @@ -105,6 +139,31 @@ void test_build_and_run() { test_1(q, k); } +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 = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, BadSource); + try { + exe_kb kbExe = syclex::build(kbSrc); + assert(false && "we should not be here"); + } catch (sycl::exception &e) { + // yas! + assert(e.code() == sycl::errc::build); + } +} + int main() { // TODO - awaiting guidance // #ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_SYCL @@ -113,7 +172,7 @@ int main() { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER test_build_and_run(); - // test_error(); + test_error(); #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif From b21ee0667308587733d240768df49a0c80b84ab6 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 16:48:00 -0700 Subject: [PATCH 08/27] implemented sycl_compilation_available --- .../kernel_compiler/kernel_compiler_sycl.cpp | 26 ++++++++++++++----- .../KernelCompiler/kernel_compiler_sycl.cpp | 4 --- 2 files changed, 19 insertions(+), 11 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 734d6aefd3200..5ca99f8dab943 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -110,6 +110,15 @@ void output_include_files(const std::filesystem::path &dpath, } } +std::string get_compiler_name() { +#ifdef __WIN32 + std::string compiler = "clang++.exe"; +#else + std::string compiler = "clang++"; +#endif + return compiler; +} + void invoke_compiler(const std::filesystem::path &fpath, const std::filesystem::path &dpath, const std::string &id, const std::vector &UserArgs, @@ -119,11 +128,7 @@ void invoke_compiler(const std::filesystem::path &fpath, std::filesystem::path parent_dir(dpath); std::filesystem::path target_path = parent_dir / (id + ".bin"); std::filesystem::path log_path = parent_dir / "compilation_log.txt"; -#ifdef __WIN32 - std::string compiler = "clang++.exe"; -#else - std::string compiler = "clang++"; -#endif + std::string compiler = get_compiler_name(); std::string command = compiler + " -fsycl -o " + target_path.make_preferred().string() + " " + @@ -203,9 +208,16 @@ spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, } bool SYCL_Compilation_Available() { - // check for clang++ clang++.exe icpx icpx.exe on PATH + // is compiler on $PATH ? We try to invoke it. + std::string id = generate_semi_unique_id(); + const std::filesystem::path tmp = std::filesystem::temp_directory_path(); + std::filesystem::path dump_path = tmp / (id + "_version.txt"); + std::string compiler = get_compiler_name(); + std::string test_command = + compiler + " --version &> " + dump_path.make_preferred().string(); + int result = std::system(test_command.c_str()); - return true; + return (result == 0); } } // namespace detail diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 950ea38d24d6a..680b102a2336d 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -165,10 +165,6 @@ void test_error() { } int main() { - // TODO - awaiting guidance - // #ifndef SYCL_EXT_ONEAPI_KERNEL_COMPILER_SYCL - // static_assert(false, "KernelCompiler SYCL feature test macro undefined"); - // #endif #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER test_build_and_run(); From 1bff1f6834eaa3af8002494fae866a95949c39ed Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 17:48:28 -0700 Subject: [PATCH 09/27] move is_source_kernel_bundle_supported out of API into detail. minor cleanup --- sycl/include/sycl/kernel_bundle.hpp | 7 ++----- sycl/source/detail/device_impl.cpp | 3 ++- sycl/source/kernel_bundle.cpp | 5 +++-- .../KernelCompiler/kernel_compiler_opencl.cpp | 13 +++++++------ .../KernelCompiler/kernel_compiler_spirv.cpp | 4 +--- .../KernelCompiler/kernel_compiler_sycl.cpp | 15 +++------------ 6 files changed, 18 insertions(+), 29 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 70e3794e5d5f6..c12c6eab3f1e7 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -865,14 +865,11 @@ struct is_property_key_of> : std::true_type {}; -///////////////////////// -// syclex::is_source_kernel_bundle_supported -///////////////////////// +namespace detail { +// forward decls __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, source_language Language); -namespace detail { -// forward decls __SYCL_EXPORT kernel_bundle make_kernel_bundle_from_source( const context &SyclContext, source_language Language, diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 6d2a8d08736f7..78dd6e6cd1136 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -857,7 +857,8 @@ bool device_impl::isGetDeviceAndHostTimerSupported() { bool device_impl::extOneapiCanCompile( ext::oneapi::experimental::source_language Language) { try { - return is_source_kernel_bundle_supported(getBackend(), Language); + return sycl::ext::oneapi::experimental::detail:: + is_source_kernel_bundle_supported(getBackend(), Language); } catch (sycl::exception &) { return false; } diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index cd35f3dc0d084..dfb80eebfa685 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -368,8 +368,10 @@ using source_kb = kernel_bundle; using exe_kb = kernel_bundle; using kernel_bundle_impl = sycl::detail::kernel_bundle_impl; +namespace detail { + ///////////////////////// -// syclex::is_source_kernel_bundle_supported +// syclex::detail::is_source_kernel_bundle_supported ///////////////////////// bool is_source_kernel_bundle_supported(backend BE, source_language Language) { // Support is limited to the opencl and level_zero backends. @@ -392,7 +394,6 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { ///////////////////////// // syclex::detail::create_kernel_bundle_from_source ///////////////////////// -namespace detail { using include_pairs_t = std::vector>; diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp index 9427442fa7efc..79b72ee19b0ae 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp @@ -85,12 +85,13 @@ void test_build_and_run() { sycl::context ctx{d}; sycl::queue q{ctx, d}; - bool ok = syclex::is_source_kernel_bundle_supported( - ctx.get_backend(), syclex::source_language::opencl); + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::opencl); if (!ok) { - std::cout << "Apparently this backend does not support OpenCL C source " + std::cout << "Apparently this device does not support OpenCL C source " "kernel bundle extension: " - << ctx.get_backend() << std::endl; + << q.get_device().get_info() + << std::endl; return; } @@ -141,8 +142,8 @@ void test_error() { sycl::context ctx{d}; sycl::queue q{ctx, d}; - bool ok = syclex::is_source_kernel_bundle_supported( - ctx.get_backend(), syclex::source_language::opencl); + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::opencl); if (!ok) { return; } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp index c0e8a7dda85ae..38567fe6ee0b3 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp @@ -175,9 +175,7 @@ void testKernelsFromSpvFile(std::string kernels_file, return bundle.ext_oneapi_get_kernel(name); }; - sycl::device d; - sycl::context ctx{d}; - sycl::queue q{ctx, d}; + sycl::queue q; auto bundle = loadKernelsFromFile(q, kernels_file); // Test simple kernel. diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 680b102a2336d..7f2b361e94021 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -80,25 +80,16 @@ void test_build_and_run() { using source_kb = sycl::kernel_bundle; using exe_kb = sycl::kernel_bundle; - // TODO: remove this dance from other tests - // this dance avoids a bug on L0, ensuring context is of exactly one device - // sycl::device d; - // sycl::context ctx{d}; - // sycl::queue q{ctx, d}; - sycl::queue q; sycl::context ctx = q.get_context(); - // TODO: replace is_source_kernel_bundle_supported() with - // device::ext_oneapi_can_compile() - // bool ok = syclex::is_source_kernel_bundle_supported(ctx.get_backend(), - // syclex::source_language::sycl); bool ok = q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); if (!ok) { - std::cout << "Apparently this backend does not support SYCL source " + std::cout << "Apparently this device does not support SYCL source " "kernel bundle extension: " - << ctx.get_backend() << std::endl; + << q.get_device().get_info() + << std::endl; return; } From dcb54ce6a46ac43cec28a3e8f8ba1cbed08b2a27 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 10 Apr 2024 18:30:04 -0700 Subject: [PATCH 10/27] win --- sycl/include/sycl/kernel_bundle.hpp | 16 ++++++---------- .../kernel_compiler/kernel_compiler_sycl.hpp | 3 ++- 2 files changed, 8 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c12c6eab3f1e7..22cb80e6d02ca 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -899,11 +899,9 @@ template , PropertyListT>::value>> -__SYCL_EXPORT kernel_bundle -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::string &Source, - PropertyListT props = {}) { +kernel_bundle create_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::string &Source, PropertyListT props = {}) { std::vector> IncludePairsVec; if constexpr (props.template has_property()) { IncludePairsVec = props.template get_property().record; @@ -920,11 +918,9 @@ template , PropertyListT>::value>> -__SYCL_EXPORT kernel_bundle -create_kernel_bundle_from_source(const context &SyclContext, - source_language Language, - const std::vector &Bytes, - PropertyListT props = {}) { +kernel_bundle create_kernel_bundle_from_source( + const context &SyclContext, source_language Language, + const std::vector &Bytes, PropertyListT props = {}) { std::vector> IncludePairsVec; if constexpr (props.template has_property()) { IncludePairsVec = props.template get_property().record; diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 799f493787a24..794aceb889fb2 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -9,9 +9,10 @@ #pragma once #include -#include // for __SYCL_EXPORT +#include // __SYCL_EXPORT #include +#include // std::accumulate #include #include From 16e3d694e55091d94642496eaa386186eb96a991 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 20 May 2024 14:41:15 -0700 Subject: [PATCH 11/27] interim registered_kernel_names implementation. this is NOT final --- .../sycl/ext/oneapi/properties/property.hpp | 3 +- sycl/include/sycl/kernel_bundle.hpp | 28 +++++++++++++++-- sycl/source/detail/kernel_bundle_impl.hpp | 6 ++-- .../kernel_compiler/kernel_compiler_sycl.cpp | 31 ++++++++++++------- .../kernel_compiler/kernel_compiler_sycl.hpp | 8 ++--- sycl/source/kernel_bundle.cpp | 13 ++++---- 6 files changed, 62 insertions(+), 27 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 081cb419ce000..1ed63af245e6c 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -208,8 +208,9 @@ enum PropKind : uint32_t { IndirectlyCallable = 67, CallsIndirectly = 68, IncludeFiles = 69, + RegisteredKernelNames = 70, // PropKindSize must always be the last value. - PropKindSize = 70, + PropKindSize = 71, }; struct property_key_base_tag {}; diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 22cb80e6d02ca..f6abb1485d148 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -865,6 +865,23 @@ struct is_property_key_of> : std::true_type {}; +///////////////////////// +// PropertyT syclex::registered_kernel_names +///////////////////////// +struct registered_kernel_names + : detail::run_time_property_key { + std::vector kernel_names; + registered_kernel_names(const std::string &knArg) : kernel_names{knArg} {} + registered_kernel_names(const std::vector &knsArg) + : kernel_names(knsArg) {} +}; +using registered_kernel_names_key = registered_kernel_names; + +template <> +struct is_property_key_of> + : std::true_type{}; + namespace detail { // forward decls __SYCL_EXPORT bool is_source_kernel_bundle_supported(backend BE, @@ -886,7 +903,8 @@ __SYCL_EXPORT kernel_bundle build_from_source(kernel_bundle &SourceKB, const std::vector &Devices, const std::vector &BuildOptions, - std::string *LogPtr); + std::string *LogPtr, + const std::vector &RegisteredKernelNames); } // namespace detail @@ -947,13 +965,19 @@ build(kernel_bundle &SourceKB, const std::vector &Devices, PropertyListT props = {}) { std::vector BuildOptionsVec; std::string *LogPtr = nullptr; + std::vector RegisteredKernelNamesVec; if constexpr (props.template has_property()) { BuildOptionsVec = props.template get_property().opts; } if constexpr (props.template has_property()) { LogPtr = props.template get_property().log; } - return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr); + if constexpr (props.template has_property()) { + RegisteredKernelNamesVec = + props.template get_property().kernel_names; + } + return detail::build_from_source(SourceKB, Devices, BuildOptionsVec, LogPtr, + RegisteredKernelNamesVec); } template build_from_source(const std::vector Devices, const std::vector &BuildOptions, - std::string *LogPtr) { + std::string *LogPtr, + const std::vector &RegisteredKernelNames) { assert(MState == bundle_state::ext_oneapi_source && "bundle_state::ext_oneapi_source required"); @@ -404,7 +405,8 @@ class kernel_bundle_impl { if (Language == syclex::source_language::sycl) { const auto &SourceStr = std::get(this->Source); return syclex::detail::SYCL_to_SPIRV(SourceStr, IncludePairs, - BuildOptions, LogPtr); + BuildOptions, LogPtr, + RegisteredKernelNames); } throw sycl::exception( make_error_code(errc::invalid), diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 5ca99f8dab943..e39b8fbdf7f36 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -65,14 +65,16 @@ void output_preamble(std::ofstream &os, const std::filesystem::path &file_path, const std::vector &UserArgs) { os << "/*\n clang++ -fsycl -o " << id << ".bin " - << user_args_as_string(UserArgs) << " -fsycl-dump-device-code=./ " << id + << user_args_as_string(UserArgs) + << "-fno-sycl-dead-args-optimization -fsycl-dump-device-code=./ " << id << ".cpp \n */" << std::endl; } -std::filesystem::path output_cpp(const std::filesystem::path &parent_dir, - const std::string &id, - std::string raw_code_string, - const std::vector &UserArgs) { +std::filesystem::path +output_cpp(const std::filesystem::path &parent_dir, const std::string &id, + std::string raw_code_string, + const std::vector &UserArgs, + const std::vector &RegisteredKernelNames) { std::filesystem::path file_path = parent_dir / (id + ".cpp"); std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); @@ -81,7 +83,11 @@ std::filesystem::path output_cpp(const std::filesystem::path &parent_dir, outfile << raw_code_string << std::endl; // temporarily needed until -c works with -fsycl-dump-spirv - outfile << "int main(){ return 0; }" << std::endl; + outfile << "int main(){\n"; + for (std::string nm : RegisteredKernelNames) { + outfile << " " << nm << ";\n"; + } + outfile << " return 0;\n}" << std::endl; outfile.close(); // Close the file when finished } else { @@ -133,7 +139,8 @@ void invoke_compiler(const std::filesystem::path &fpath, std::string command = compiler + " -fsycl -o " + target_path.make_preferred().string() + " " + user_args_as_string(UserArgs) + - " -fsycl-dump-device-code=" + parent_dir.make_preferred().string() + " " + + " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=" + + parent_dir.make_preferred().string() + " " + file_path.make_preferred().string() + " 2> " + log_path.make_preferred().string(); @@ -192,14 +199,14 @@ spirv_vec_t load_spv_from_file(std::filesystem::path file_name) { return spv; } -spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, - include_pairs_t IncludePairs, - const std::vector &UserArgs, - std::string *LogPtr) { +spirv_vec_t +SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames) { // clang-format off const std::string id = generate_semi_unique_id(); const std::filesystem::path parent_dir = prepare_ws(id); - std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource, UserArgs); + std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource, UserArgs, RegisteredKernelNames); output_include_files(parent_dir, IncludePairs); invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); std::filesystem::path spv_path = find_spv(parent_dir, id); diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 794aceb889fb2..dfff9ac839e84 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -24,10 +24,10 @@ namespace detail { using spirv_vec_t = std::vector; using include_pairs_t = std::vector>; -spirv_vec_t SYCL_to_SPIRV(const std::string &Source, - include_pairs_t IncludePairs, - const std::vector &UserArgs, - std::string *LogPtr); +spirv_vec_t +SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames); bool SYCL_Compilation_Available(); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index dfb80eebfa685..4acc60fedaef2 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -439,15 +439,16 @@ source_kb make_kernel_bundle_from_source(const context &SyclContext, // syclex::detail::build_from_source(source_kb) => exe_kb ///////////////////////// -exe_kb build_from_source(source_kb &SourceKB, - const std::vector &Devices, - const std::vector &BuildOptions, - std::string *LogPtr) { +exe_kb +build_from_source(source_kb &SourceKB, const std::vector &Devices, + const std::vector &BuildOptions, + std::string *LogPtr, + const std::vector &RegisteredKernelNames) { std::vector UniqueDevices = sycl::detail::removeDuplicateDevices(Devices); std::shared_ptr sourceImpl = getSyclObjImpl(SourceKB); - std::shared_ptr KBImpl = - sourceImpl->build_from_source(UniqueDevices, BuildOptions, LogPtr); + std::shared_ptr KBImpl = sourceImpl->build_from_source( + UniqueDevices, BuildOptions, LogPtr, RegisteredKernelNames); return sycl::detail::createSyclObjFromImpl(KBImpl); } From 7163a1d2a272dde97368959315eba11a705ecbb5 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 13 Jun 2024 14:22:30 -0700 Subject: [PATCH 12/27] merge conflics and test bump --- sycl/include/sycl/kernel_bundle.hpp | 6 +++--- .../KernelCompiler/kernel_compiler_sycl.cpp | 13 ++++++------- 2 files changed, 9 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index f6abb1485d148..56a8f0450f8dd 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -911,7 +911,7 @@ build_from_source(kernel_bundle &SourceKB, ///////////////////////// // syclex::create_kernel_bundle_from_source ///////////////////////// -template && detail::all_props_are_keys_of< @@ -930,7 +930,7 @@ kernel_bundle create_kernel_bundle_from_source( } #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) -template && detail::all_props_are_keys_of< @@ -953,7 +953,7 @@ kernel_bundle create_kernel_bundle_from_source( // syclex::build(source_kb) => exe_kb ///////////////////////// -template && detail::all_props_are_keys_of< diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 7f2b361e94021..eb85b490fe2cb 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -24,7 +24,7 @@ auto constexpr SYCLSource = R"===( #include #include "AddEm.h" -SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) void ff_cp(int *ptr) { // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>() @@ -38,7 +38,7 @@ void ff_cp(int *ptr) { auto constexpr BadSource = R"===( #include -SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) void ff_cp(int *ptr) { sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); @@ -117,12 +117,11 @@ void test_build_and_run() { assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); - // amend __free_function_ to kernel f name. - sycl::kernel k = kbExe2.ext_oneapi_get_kernel("__free_function_ff_cp"); + sycl::kernel k = kbExe2.ext_oneapi_get_kernel("_Z19__sycl_kernel_ff_cp"); - // NOTE THIS NOISE - // sycl::kernel_bundle kb = - // syclexp::build(kb_src, + // COMING SOON + // sycl::kernel_bundle kb + // = syclexp::build(kb_src, // syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); // sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar"); From 3c048c3764ae995aed17d243fb72a4c5f32b7206 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 19 Jun 2024 10:13:51 -0700 Subject: [PATCH 13/27] clang-format and linux symbols --- sycl/include/sycl/kernel_bundle.hpp | 2 +- sycl/test/abi/sycl_symbols_linux.dump | 24 ++++++++++++------------ 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 56a8f0450f8dd..ed48f730911a4 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -880,7 +880,7 @@ using registered_kernel_names_key = registered_kernel_names; template <> struct is_property_key_of> - : std::true_type{}; + : std::true_type {}; namespace detail { // forward decls diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 0edaaa25b4ba1..32c73732dae43 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3050,9 +3050,6 @@ _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_V13ext6oneapi12experimental32create_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaIS9_EE -_ZN4sycl3_V13ext6oneapi12experimental33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi1EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi2EEEvNS0_5rangeIXT_EEE _ZN4sycl3_V13ext6oneapi12experimental4node12update_rangeILi3EEEvNS0_5rangeIXT_EEE @@ -3064,7 +3061,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC1ERKNS3_16image_des _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_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bundleILNS0_12bundle_stateE3EEERKSt6vectorINS0_6deviceESaISA_EERKS9_INSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISK_EEPSK_SO_ _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKvm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base14updateAccessorEPKNS0_6detail16AccessorBaseHostE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv @@ -3088,6 +3085,9 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_5 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_5queueERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESt6vectorISt4pairISE_SE_ESaISJ_EE +_ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaISA_EES9_ISt4pairINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESL_ESaISM_EE +_ZN4sycl3_V13ext6oneapi12experimental6detail33is_source_kernel_bundle_supportedENS0_7backendENS3_15source_languageE _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental9image_memC2ERKNS3_16image_descriptorERKNS0_5queueE @@ -3119,6 +3119,7 @@ _ZN4sycl3_V15queue10mem_adviseEPKvmiRKSt6vectorINS0_5eventESaIS5_EERKNS0_6detail _ZN4sycl3_V15queue10wait_proxyERKNS0_6detail13code_locationE _ZN4sycl3_V15queue11submit_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationE _ZN4sycl3_V15queue11submit_implESt8functionIFvRNS0_7handlerEEES1_RKNS0_6detail13code_locationE +_ZN4sycl3_V15queue15ext_oneapi_prodEv _ZN4sycl3_V15queue17discard_or_returnERKNS0_5eventE _ZN4sycl3_V15queue18throw_asynchronousEv _ZN4sycl3_V15queue20memcpyToDeviceGlobalEPvPKvbmmRKSt6vectorINS0_5eventESaIS6_EE @@ -3621,6 +3622,7 @@ _ZN4sycl3_V17handler28memcpyToHostOnlyDeviceGlobalEPKvS3_mbmm _ZN4sycl3_V17handler28setStateExplicitKernelBundleEv _ZN4sycl3_V17handler30memcpyFromHostOnlyDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler30verifyUsedKernelBundleInternalENS0_6detail11string_viewE +_ZN4sycl3_V17handler32verifyDeviceHasProgressGuaranteeENS0_3ext6oneapi12experimental26forward_progress_guaranteeENS4_15execution_scopeES6_ _ZN4sycl3_V17handler34ext_oneapi_wait_external_semaphoreENS0_3ext6oneapi12experimental24interop_semaphore_handleE _ZN4sycl3_V17handler36ext_oneapi_signal_external_semaphoreENS0_3ext6oneapi12experimental24interop_semaphore_handleE _ZN4sycl3_V17handler6memcpyEPvPKvm @@ -3633,7 +3635,6 @@ _ZN4sycl3_V17handlerC1ESt10shared_ptrINS0_6detail10queue_implEEb _ZN4sycl3_V17handlerC2ESt10shared_ptrINS0_3ext6oneapi12experimental6detail10graph_implEE _ZN4sycl3_V17handlerC2ESt10shared_ptrINS0_6detail10queue_implEES5_S5_b _ZN4sycl3_V17handlerC2ESt10shared_ptrINS0_6detail10queue_implEEb -_ZN4sycl3_V17handler32verifyDeviceHasProgressGuaranteeENS0_3ext6oneapi12experimental26forward_progress_guaranteeENS4_15execution_scopeES6_ _ZN4sycl3_V17samplerC1ENS0_29coordinate_normalization_modeENS0_15addressing_modeENS0_14filtering_modeERKNS0_13property_listE _ZN4sycl3_V17samplerC1EP11_cl_samplerRKNS0_7contextE _ZN4sycl3_V17samplerC2ENS0_29coordinate_normalization_modeENS0_15addressing_modeENS0_14filtering_modeERKNS0_13property_listE @@ -3748,7 +3749,6 @@ _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue16enable_profilingEEEbv _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue4cuda18use_default_streamEEEbv _ZNK4sycl3_V15queue12has_propertyINS0_8property5queue8in_orderEEEbv _ZNK4sycl3_V15queue16ext_oneapi_emptyEv -_ZN4sycl3_V15queue15ext_oneapi_prodEv _ZNK4sycl3_V15queue16get_backend_infoINS0_4info6device15backend_versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv _ZNK4sycl3_V15queue16get_backend_infoINS0_4info6device7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv _ZNK4sycl3_V15queue16get_backend_infoINS0_4info8platform7versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv @@ -3973,6 +3973,12 @@ _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device22m _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device22max_image_linear_widthEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device23max_image_linear_heightEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device26max_image_linear_row_pitchEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31sub_group_progress_capabilitiesILNS5_15execution_scopeE2EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31sub_group_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE1EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE2EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv +_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device32work_group_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_3ext8codeplay12experimental4info6device15supports_fusionEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_3ext8codeplay12experimental4info6device28max_registers_per_work_groupEEENS0_6detail11ABINeutralTINS9_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_4info6device10extensionsEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv @@ -4084,12 +4090,6 @@ _ZNK4sycl3_V16device13get_info_implINS0_4info6device7versionEEENS0_6detail11ABIN _ZNK4sycl3_V16device13get_info_implINS0_4info6device8atomic64EEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_4info6device8platformEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13get_info_implINS0_4info6device9vendor_idEEENS0_6detail11ABINeutralTINS6_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device32work_group_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31sub_group_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31sub_group_progress_capabilitiesILNS5_15execution_scopeE2EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE2EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE3EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv -_ZNK4sycl3_V16device13get_info_implINS0_3ext6oneapi12experimental4info6device31work_item_progress_capabilitiesILNS5_15execution_scopeE1EEEEENS0_6detail11ABINeutralTINSB_19is_device_info_descIT_E11return_typeEE4typeEv _ZNK4sycl3_V16device13has_extensionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK4sycl3_V16device14is_acceleratorEv _ZNK4sycl3_V16device16get_backend_infoINS0_4info6device15backend_versionEEENS0_6detail20is_backend_info_descIT_E11return_typeEv From 8dbc79b2cf0b75eb1bad01958e187e27e85e2a61 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 19 Jun 2024 15:01:20 -0700 Subject: [PATCH 14/27] test updates --- sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp | 4 ++-- .../EnqueueFunctions/kernel_submit_with_event_and_kb.cpp | 4 ++-- sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp | 4 ++-- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp | 3 ++- 4 files changed, 8 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp index 7176def8ec7de..485266c8a3d06 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_shortcut_with_kb.cpp @@ -21,8 +21,8 @@ constexpr size_t N = 1024; int main() { sycl::queue Q; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp index 0e8988574fd5d..b2731740a60e3 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_event_and_kb.cpp @@ -21,8 +21,8 @@ int main() { sycl::queue Q; int Memory[N] = {0}; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp index be29a73b87ee0..2651a803b509c 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_submit_with_kb.cpp @@ -21,8 +21,8 @@ int main() { sycl::queue Q; int Memory[N] = {0}; - if (!oneapiext::is_source_kernel_bundle_supported( - Q.get_backend(), oneapiext::source_language::opencl)) { + if (!Q.get_device().ext_oneapi_can_compile( + oneapiext::source_language::opencl)) { std::cout << "Backend does not support OpenCL C source kernel bundle extension: " << Q.get_backend() << std::endl; diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index eb85b490fe2cb..58586ea2ab978 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -12,7 +12,8 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include +#include +#include auto constexpr AddEmH = R"===( int AddEm(int a, int b){ From ae8f9366ab3c20798d2e98de19a286723e030251 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 20 Jun 2024 08:38:44 -0700 Subject: [PATCH 15/27] resolve merge conflicts and bump counter for the #include that is in the string of the kernel source --- sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp index 7b2aa6b3b5883..74f8e1aed8aff 100644 --- a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp @@ -7,7 +7,7 @@ // CHECK-DAG: no_sycl_hpp_in_e2e_tests.cpp // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 3 +// CHECK-NUM-MATCHES: 4 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see From 531bcbc4b2a8fc3698564ae5f899db698d6dba0e Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Thu, 20 Jun 2024 09:47:46 -0700 Subject: [PATCH 16/27] win fix and symbols --- .../kernel_compiler/kernel_compiler_sycl.cpp | 3 +- sycl/test/abi/sycl_symbols_windows.dump | 28 +++++++++---------- 2 files changed, 16 insertions(+), 15 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index e39b8fbdf7f36..c751c392be175 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include "kernel_compiler_sycl.hpp" @@ -230,4 +231,4 @@ bool SYCL_Compilation_Available() { } // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 -} // namespace sycl \ No newline at end of file +} // namespace sycl diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index e8610211e8572..2240cef7cddb6 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -41,18 +41,12 @@ ??$get_info@U?$max_work_groups@$00@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$id@$00@23@XZ ??$get_info@U?$max_work_groups@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$id@$01@23@XZ ??$get_info@U?$max_work_groups@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$id@$02@23@XZ +??$get_info@U?$sub_group_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@U?$sub_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info@U?$work_item_progress_capabilities@$00@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@U?$work_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$work_item_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info@U?$work_item_progress_capabilities@$00@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@U?$work_item_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$work_item_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@U?$work_item_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$sub_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$work_item_progress_capabilities@$00@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info@U?$sub_group_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$work_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ -??$get_info_impl@U?$sub_group_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info@Uarchitecture@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA?AW4architecture@experimental@oneapi@ext@23@XZ ??$get_info@Uatomic_fence_order_capabilities@context@info@_V1@sycl@@@context@_V1@sycl@@QEBA?AV?$vector@W4memory_order@_V1@sycl@@V?$allocator@W4memory_order@_V1@sycl@@@std@@@std@@XZ ??$get_info@Uatomic_fence_scope_capabilities@context@info@_V1@sycl@@@context@_V1@sycl@@QEBA?AV?$vector@W4memory_scope@_V1@sycl@@V?$allocator@W4memory_scope@_V1@sycl@@@std@@@std@@XZ @@ -108,6 +102,12 @@ ??$get_info_impl@U?$max_work_item_sizes@$00@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$range@$00@12@XZ ??$get_info_impl@U?$max_work_item_sizes@$01@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$range@$01@12@XZ ??$get_info_impl@U?$max_work_item_sizes@$02@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$range@$02@12@XZ +??$get_info_impl@U?$sub_group_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@U?$sub_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@U?$work_group_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@U?$work_item_progress_capabilities@$00@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@U?$work_item_progress_capabilities@$01@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ +??$get_info_impl@U?$work_item_progress_capabilities@$02@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4forward_progress_guarantee@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ ??$get_info_impl@Uaddress_bits@device@info@_V1@sycl@@@device@_V1@sycl@@AEBAIXZ ??$get_info_impl@Uarchitecture@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@AEBA?AW4architecture@experimental@oneapi@ext@12@XZ ??$get_info_impl@Uaspects@device@info@_V1@sycl@@@device@_V1@sycl@@AEBA?AV?$vector@W4aspect@_V1@sycl@@V?$allocator@W4aspect@_V1@sycl@@@std@@@std@@XZ @@ -4003,7 +4003,7 @@ ?begin@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAXAEAVqueue@67@AEBVproperty_list@67@@Z ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAXAEBV?$vector@Vqueue@_V1@sycl@@V?$allocator@Vqueue@_V1@sycl@@@std@@@std@@AEBVproperty_list@67@@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_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@@2@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 @@ -4046,8 +4046,6 @@ ?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 -?create_kernel_bundle_from_source@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@45@AEBVcontext@45@W4source_language@12345@AEBV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@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 @@ -4080,7 +4078,6 @@ ?ext_intel_read_host_pipe@handler@_V1@sycl@@AEAAXVstring_view@detail@23@PEAX_K_N@Z ?ext_intel_write_host_pipe@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEAX_K_N@Z ?ext_intel_write_host_pipe@handler@_V1@sycl@@AEAAXVstring_view@detail@23@PEAX_K_N@Z -?verifyDeviceHasProgressGuarantee@handler@_V1@sycl@@AEAAXW4forward_progress_guarantee@experimental@oneapi@ext@23@W4execution_scope@56723@1@Z ?ext_oneapi_advise_usm_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEBX_KW4_pi_mem_advice@@V?$vector@IV?$allocator@I@std@@@6@PEAI@Z ?ext_oneapi_architecture_is@device@_V1@sycl@@QEAA_NW4arch_category@experimental@oneapi@ext@23@@Z ?ext_oneapi_architecture_is@device@_V1@sycl@@QEAA_NW4architecture@experimental@oneapi@ext@23@@Z @@ -4096,7 +4093,6 @@ ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@0AEBUimage_descriptor@56723@@Z ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@56723@@Z ?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@56723@PEAX111@Z -?ext_oneapi_prod@queue@_V1@sycl@@QEAAXXZ ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBUcode_location@detail@23@@Z ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z ?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KV423@AEBUcode_location@detail@23@@Z @@ -4158,6 +4154,7 @@ ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vstream@_V1@sycl@@@2oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVstream@34@@Z ?ext_oneapi_prefetch_usm_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAX_KV?$vector@IV?$allocator@I@std@@@6@PEAI@Z +?ext_oneapi_prod@queue@_V1@sycl@@QEAAXXZ ?ext_oneapi_set_external_event@queue@_V1@sycl@@QEAAXAEBVevent@23@@Z ?ext_oneapi_signal_external_semaphore@handler@_V1@sycl@@QEAAXUinterop_semaphore_handle@experimental@oneapi@ext@23@@Z ?ext_oneapi_signal_external_semaphore@handler@_V1@sycl@@QEAAXUinterop_semaphore_handle@experimental@oneapi@ext@23@_K@Z @@ -4455,7 +4452,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_source_kernel_bundle_supported@detail@experimental@oneapi@ext@_V1@sycl@@YA_NW4backend@56@W4source_language@23456@@Z ?is_specialization_constant_set@kernel_bundle_plain@detail@_V1@sycl@@IEBA_NPEBD@Z ?join_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$vector@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@2@@5@W4bundle_state@23@@Z ?lgamma_r_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@PEAH@Z @@ -4478,6 +4475,8 @@ ?make_kernel@detail@_V1@sycl@@YA?AVkernel@23@_KAEBVcontext@23@W4backend@23@@Z ?make_kernel_bundle@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@_KAEBVcontext@23@W4bundle_state@23@W4backend@23@@Z ?make_kernel_bundle@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@_KAEBVcontext@23@_NW4bundle_state@23@W4backend@23@@Z +?make_kernel_bundle_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@56@AEBVcontext@56@W4source_language@23456@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$vector@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@V?$allocator@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@@2@@std@@@Z +?make_kernel_bundle_from_source@detail@experimental@oneapi@ext@_V1@sycl@@YA?AV?$kernel_bundle@$02@56@AEBVcontext@56@W4source_language@23456@AEBV?$vector@W4byte@std@@V?$allocator@W4byte@std@@@2@@std@@V?$vector@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@V?$allocator@U?$pair@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V12@@std@@@2@@std@@@Z ?make_platform@detail@_V1@sycl@@YA?AVplatform@23@_KW4backend@23@@Z ?make_platform@level_zero@oneapi@ext@_V1@sycl@@YA?AVplatform@45@_K@Z ?make_platform@opencl@_V1@sycl@@YA?AVplatform@23@_K@Z @@ -4655,6 +4654,7 @@ ?useHostPtr@SYCLMemObjT@detail@_V1@sycl@@QEAA_NXZ ?use_kernel_bundle@handler@_V1@sycl@@QEAAXAEBV?$kernel_bundle@$01@23@@Z ?usesPinnedHostMemory@SYCLMemObjT@detail@_V1@sycl@@UEBA_NXZ +?verifyDeviceHasProgressGuarantee@handler@_V1@sycl@@AEAAXW4forward_progress_guarantee@experimental@oneapi@ext@23@W4execution_scope@56723@1@Z ?verifyKernelInvoc@handler@_V1@sycl@@AEAAXAEBVkernel@23@@Z ?verifyUsedKernelBundle@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z ?verifyUsedKernelBundleInternal@handler@_V1@sycl@@AEAAXVstring_view@detail@23@@Z From e3cc1d274884f9318af7d8794f0fe1a552dada75 Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Thu, 20 Jun 2024 10:35:30 -0700 Subject: [PATCH 17/27] std::byte screen --- sycl/include/sycl/kernel_bundle.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index ed48f730911a4..c9a9175d42504 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -893,11 +893,13 @@ make_kernel_bundle_from_source( const std::string &Source, std::vector> IncludePairsVec); +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) __SYCL_EXPORT kernel_bundle make_kernel_bundle_from_source( const context &SyclContext, source_language Language, const std::vector &Bytes, std::vector> IncludePairsVec); +#endif __SYCL_EXPORT kernel_bundle build_from_source(kernel_bundle &SourceKB, From 4445525178041cd3f44f22f23561ab03247cb7bf Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 24 Jun 2024 09:23:31 -0700 Subject: [PATCH 18/27] patch from rajiv for name mangling fix on extern C and concominant test update --- clang/lib/Sema/SemaSYCL.cpp | 26 ++++++++++++------- .../KernelCompiler/kernel_compiler_sycl.cpp | 2 +- 2 files changed, 17 insertions(+), 11 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7a0b2dede0a7f..27bdf6966e9c1 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1111,18 +1111,24 @@ static std::pair constructFreeFunctionKernelName( SemaSYCL &SemaSYCLRef, const FunctionDecl *FreeFunc, MangleContext &MC) { SmallString<256> Result; llvm::raw_svector_ostream Out(Result); + std::string NewName; std::string StableName; - MC.mangleName(FreeFunc, Out); - std::string MangledName(Out.str()); - size_t StartNums = MangledName.find_first_of("0123456789"); - size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums); - size_t NameLength = - std::stoi(MangledName.substr(StartNums, EndNums - StartNums)); - size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength; - std::string NewName = MangledName.substr(0, StartNums) + - std::to_string(NewNameLength) + "__sycl_kernel_" + - MangledName.substr(EndNums); + // Handle extern "C" + if (FreeFunc->getLanguageLinkage() == CLanguageLinkage) { + const IdentifierInfo *II = FreeFunc->getIdentifier(); + NewName = "__sycl_kernel_" + II->getName().str(); + } else { + MC.mangleName(FreeFunc, Out); + std::string MangledName(Out.str()); + size_t StartNums = MangledName.find_first_of("0123456789"); + size_t EndNums = MangledName.find_first_not_of("0123456789", StartNums); + size_t NameLength = + std::stoi(MangledName.substr(StartNums, EndNums - StartNums)); + size_t NewNameLength = 14 /*length of __sycl_kernel_*/ + NameLength; + NewName = MangledName.substr(0, StartNums) + std::to_string(NewNameLength) + + "__sycl_kernel_" + MangledName.substr(EndNums); + } StableName = NewName; return {NewName, StableName}; } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 58586ea2ab978..681b47fdeea74 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -118,7 +118,7 @@ void test_build_and_run() { assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); - sycl::kernel k = kbExe2.ext_oneapi_get_kernel("_Z19__sycl_kernel_ff_cp"); + sycl::kernel k = kbExe2.ext_oneapi_get_kernel("__sycl_kernel_ff_cp"); // COMING SOON // sycl::kernel_bundle kb From 424da1b162f1994a080b17c7156194836387d6be Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 24 Jun 2024 15:44:10 -0700 Subject: [PATCH 19/27] excluding kernel_compiler+sycl from GCC<8 b.c. std:filesystem. --- .../kernel_compiler/kernel_compiler_sycl.cpp | 20 ++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index c751c392be175..3029d3f7c3a47 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -7,6 +7,23 @@ //===----------------------------------------------------------------------===// #include // make_error_code +#include "kernel_compiler_sycl.hpp" + +#if __GNUC__ && __GNUC__ < 8 + // std::filesystem is not availalbe for GCC < 8 + // and much of the cross-platform file handling code depends upon it. + // Given that this extension is experimental and that the file + // handling aspects are most likely temporary, it makes sense to + // simply not support GCC<8. + bool SYCL_Compilation_Available() { return false; } + spirv_vec_t +SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames){ + throw sycl::exception(sycl::errc::build, "kernel_compiler does not supprot GCC<8"); + } + +#else #include #include @@ -15,7 +32,7 @@ #include #include -#include "kernel_compiler_sycl.hpp" + namespace sycl { inline namespace _V1 { @@ -232,3 +249,4 @@ bool SYCL_Compilation_Available() { } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl +#endif From 0c9d024495ca75a864e9a8acb7d23ae1b418616c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 24 Jun 2024 16:09:27 -0700 Subject: [PATCH 20/27] clang-format? --- .../kernel_compiler/kernel_compiler_sycl.cpp | 24 +++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 3029d3f7c3a47..31970724045bc 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -6,23 +6,23 @@ // //===----------------------------------------------------------------------===// -#include // make_error_code #include "kernel_compiler_sycl.hpp" +#include // make_error_code #if __GNUC__ && __GNUC__ < 8 - // std::filesystem is not availalbe for GCC < 8 - // and much of the cross-platform file handling code depends upon it. - // Given that this extension is experimental and that the file - // handling aspects are most likely temporary, it makes sense to - // simply not support GCC<8. - bool SYCL_Compilation_Available() { return false; } - spirv_vec_t +// std::filesystem is not availalbe for GCC < 8 +// and much of the cross-platform file handling code depends upon it. +// Given that this extension is experimental and that the file +// handling aspects are most likely temporary, it makes sense to +// simply not support GCC<8. +bool SYCL_Compilation_Available() { return false; } +spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, - const std::vector &RegisteredKernelNames){ - throw sycl::exception(sycl::errc::build, "kernel_compiler does not supprot GCC<8"); - } - + const std::vector &RegisteredKernelNames) { + throw sycl::exception(sycl::errc::build, + "kernel_compiler does not supprot GCC<8"); +} #else #include From 934d8faa7fa9e18f182248d9ca569b8b93f82022 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 26 Jun 2024 08:48:11 -0700 Subject: [PATCH 21/27] clang-format and ns fix --- .../kernel_compiler/kernel_compiler_sycl.cpp | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 31970724045bc..0ca0b313158ca 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -10,12 +10,20 @@ #include // make_error_code #if __GNUC__ && __GNUC__ < 8 + // std::filesystem is not availalbe for GCC < 8 // and much of the cross-platform file handling code depends upon it. // Given that this extension is experimental and that the file // handling aspects are most likely temporary, it makes sense to // simply not support GCC<8. + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + bool SYCL_Compilation_Available() { return false; } + spirv_vec_t SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, @@ -23,6 +31,11 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, throw sycl::exception(sycl::errc::build, "kernel_compiler does not supprot GCC<8"); } +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + #else #include @@ -32,8 +45,6 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, #include #include - - namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { From 27ad9d7792d46bd52bcefb72303c38668636d2a2 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 27 Jun 2024 13:31:41 -0700 Subject: [PATCH 22/27] reviewer feedback --- sycl/include/sycl/kernel_bundle.hpp | 51 ++++++++++--------- .../KernelCompiler/kernel_compiler_sycl.cpp | 1 + 2 files changed, 27 insertions(+), 25 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index c9a9175d42504..18ab1bfe5f994 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -814,6 +814,11 @@ build(const kernel_bundle &InputBundle, namespace ext::oneapi::experimental { +namespace detail { +struct create_bundle_from_source_props; +struct build_source_bundle_props; +} // namespace detail + ///////////////////////// // PropertyT syclex::include_files ///////////////////////// @@ -832,7 +837,7 @@ using include_files_key = include_files; template <> struct is_property_key_of> + detail::create_bundle_from_source_props> : std::true_type {}; ///////////////////////// @@ -847,8 +852,7 @@ struct build_options using build_options_key = build_options; template <> -struct is_property_key_of> +struct is_property_key_of : std::true_type {}; ///////////////////////// @@ -861,8 +865,7 @@ struct save_log : detail::run_time_property_key { using save_log_key = save_log; template <> -struct is_property_key_of> +struct is_property_key_of : std::true_type {}; ///////////////////////// @@ -879,8 +882,8 @@ using registered_kernel_names_key = registered_kernel_names; template <> struct is_property_key_of> - : std::true_type {}; + detail::build_source_bundle_props> : std::true_type { +}; namespace detail { // forward decls @@ -913,12 +916,12 @@ build_from_source(kernel_bundle &SourceKB, ///////////////////////// // syclex::create_kernel_bundle_from_source ///////////////////////// -template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> +template < + typename PropertyListT = empty_properties_t, + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_are_keys_of::value>> kernel_bundle create_kernel_bundle_from_source( const context &SyclContext, source_language Language, const std::string &Source, PropertyListT props = {}) { @@ -932,12 +935,12 @@ kernel_bundle create_kernel_bundle_from_source( } #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) -template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> +template < + typename PropertyListT = empty_properties_t, + typename = std::enable_if_t< + is_property_list_v && + detail::all_props_are_keys_of::value>> kernel_bundle create_kernel_bundle_from_source( const context &SyclContext, source_language Language, const std::vector &Bytes, PropertyListT props = {}) { @@ -958,9 +961,8 @@ kernel_bundle create_kernel_bundle_from_source( template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> + detail::all_props_are_keys_of::value>> kernel_bundle build(kernel_bundle &SourceKB, @@ -985,9 +987,8 @@ build(kernel_bundle &SourceKB, template && - detail::all_props_are_keys_of< - kernel_bundle, - PropertyListT>::value>> + detail::all_props_are_keys_of::value>> kernel_bundle build(kernel_bundle &SourceKB, PropertyListT props = {}) { diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 681b47fdeea74..7093509693a9f 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -21,6 +21,7 @@ auto constexpr AddEmH = R"===( } )==="; +// TODO: remove SYCL_EXTERNAL once it is no longer needed. auto constexpr SYCLSource = R"===( #include #include "AddEm.h" From ad634bdd772fd4747db2d01ff9a9cb99c470de2a Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 27 Jun 2024 16:16:57 -0700 Subject: [PATCH 23/27] reviewer feedback, but needs cleanup --- sycl/source/detail/kernel_bundle_impl.hpp | 40 +++++++++++++++---- .../KernelCompiler/kernel_compiler_sycl.cpp | 24 ++++++++++- 2 files changed, 55 insertions(+), 9 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 5ae51d6c26323..17ece657533b4 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -352,10 +352,14 @@ class kernel_bundle_impl { // interop constructor kernel_bundle_impl(context Ctx, std::vector Devs, device_image_plain &DevImage, - std::vector KNames) + std::vector KNames, + std::vector RegKNames, + syclex::source_language Lang) : kernel_bundle_impl(Ctx, Devs, DevImage) { MState = bundle_state::executable; KernelNames = KNames; + RegisteredKernelNamesVec = RegKNames; + Language = Lang; } std::shared_ptr @@ -447,12 +451,28 @@ class kernel_bundle_impl { nullptr, MContext, MDevices, bundle_state::executable, KernelIDs, PiProgram); device_image_plain DevImg{DevImgImpl}; - return std::make_shared(MContext, MDevices, DevImg, - KernelNames); + return std::make_shared( + MContext, MDevices, DevImg, KernelNames, RegisteredKernelNames, + Language); + } + + std::string adjust_kernel_name(const std::string &Name, + syclex::source_language Lang) { + // Once name demangling support is in, we won't need this. + if (Lang != syclex::source_language::sycl) + return Name; + + bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; + return isMangled ? Name : "__sycl_kernel_" + Name; + + // bool isRegisteredName = std::find(RegisteredKernelNamesVec.begin(), + // RegisteredKernelNamesVec.end(), Name) != RegisteredKernelNamesVec.end(); + // return isRegisteredName ? Name : "__sycl_kernel_" + Name; } bool ext_oneapi_has_kernel(const std::string &Name) { - auto it = std::find(KernelNames.begin(), KernelNames.end(), Name); + auto it = std::find(KernelNames.begin(), KernelNames.end(), + adjust_kernel_name(Name, Language)); return it != KernelNames.end(); } @@ -465,9 +485,11 @@ class kernel_bundle_impl { "kernel_bundles successfully built from " "kernel_bundle."); + std::string AdjustedName = adjust_kernel_name(Name, Language); if (!ext_oneapi_has_kernel(Name)) throw sycl::exception(make_error_code(errc::invalid), - "kernel '" + Name + "' not found in kernel_bundle"); + "kernel '" + AdjustedName + + "' not found in kernel_bundle"); assert(MDeviceImages.size() > 0); const std::shared_ptr &DeviceImageImpl = @@ -476,7 +498,8 @@ class kernel_bundle_impl { ContextImplPtr ContextImpl = getSyclObjImpl(MContext); const PluginPtr &Plugin = ContextImpl->getPlugin(); sycl::detail::pi::PiKernel PiKernel = nullptr; - Plugin->call(PiProgram, Name.c_str(), &PiKernel); + Plugin->call(PiProgram, AdjustedName.c_str(), + &PiKernel); // Kernel created by piKernelCreate is implicitly retained. std::shared_ptr KernelImpl = std::make_shared( @@ -721,11 +744,14 @@ class kernel_bundle_impl { SpecConstMapT MSpecConstValues; bool MIsInterop = false; bundle_state MState; + // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames, IncludePairs - const syclex::source_language Language = syclex::source_language::opencl; + // Language is for both state::source and state::executable + syclex::source_language Language = syclex::source_language::opencl; const std::variant> Source; // only kernel_bundles created from source have KernelNames member. std::vector KernelNames; + std::vector RegisteredKernelNamesVec; include_pairs_t IncludePairs; }; diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 7093509693a9f..8784c5dffca56 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -26,6 +26,7 @@ auto constexpr SYCLSource = R"===( #include #include "AddEm.h" +// use extern "C" to avoid name mangling extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) void ff_cp(int *ptr) { @@ -35,6 +36,17 @@ void ff_cp(int *ptr) { sycl::id<1> GId = Item.get_global_id(); ptr[GId.get(0)] = AddEm(GId.get(0), 37); } + +// this name will be mangled +template +SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_templated(T *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = GId.get(0) + 39; +} )==="; auto constexpr BadSource = R"===( @@ -115,11 +127,18 @@ void test_build_and_run() { std::vector devs = kbSrc.get_devices(); exe_kb kbExe2 = syclex::build( kbSrc, devs, - syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}}); + syclex::properties{// syclex::build_options{flags}, + syclex::save_log{&log}, + syclex::registered_kernel_names{"ff_templated"}}); assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); - sycl::kernel k = kbExe2.ext_oneapi_get_kernel("__sycl_kernel_ff_cp"); + // extern "C" was used, so the name "ff_cp" is not mangled + sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); + // the templated function name will have been mangled. Mapping from original + // name to mangled is not yet supported. + sycl::kernel k2 = + kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); // COMING SOON // sycl::kernel_bundle kb @@ -129,6 +148,7 @@ void test_build_and_run() { // 4 test_1(q, k); + test_1(q, k2); } void test_error() { From cec9a7d4c9f47f8856884bdeef1624e2bc3bbcf2 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 27 Jun 2024 16:58:05 -0700 Subject: [PATCH 24/27] clean up and test improvements --- sycl/source/detail/kernel_bundle_impl.hpp | 12 ++-------- .../KernelCompiler/kernel_compiler_sycl.cpp | 23 +++++++------------ 2 files changed, 10 insertions(+), 25 deletions(-) diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 17ece657533b4..e7052b8b7c405 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -353,12 +353,10 @@ class kernel_bundle_impl { kernel_bundle_impl(context Ctx, std::vector Devs, device_image_plain &DevImage, std::vector KNames, - std::vector RegKNames, syclex::source_language Lang) : kernel_bundle_impl(Ctx, Devs, DevImage) { MState = bundle_state::executable; KernelNames = KNames; - RegisteredKernelNamesVec = RegKNames; Language = Lang; } @@ -451,9 +449,8 @@ class kernel_bundle_impl { nullptr, MContext, MDevices, bundle_state::executable, KernelIDs, PiProgram); device_image_plain DevImg{DevImgImpl}; - return std::make_shared( - MContext, MDevices, DevImg, KernelNames, RegisteredKernelNames, - Language); + return std::make_shared(MContext, MDevices, DevImg, + KernelNames, Language); } std::string adjust_kernel_name(const std::string &Name, @@ -464,10 +461,6 @@ class kernel_bundle_impl { bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; return isMangled ? Name : "__sycl_kernel_" + Name; - - // bool isRegisteredName = std::find(RegisteredKernelNamesVec.begin(), - // RegisteredKernelNamesVec.end(), Name) != RegisteredKernelNamesVec.end(); - // return isRegisteredName ? Name : "__sycl_kernel_" + Name; } bool ext_oneapi_has_kernel(const std::string &Name) { @@ -751,7 +744,6 @@ class kernel_bundle_impl { const std::variant> Source; // only kernel_bundles created from source have KernelNames member. std::vector KernelNames; - std::vector RegisteredKernelNamesVec; include_pairs_t IncludePairs; }; diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index 8784c5dffca56..ebf0e15ce1751 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -62,7 +62,7 @@ void ff_cp(int *ptr) { } )==="; -void test_1(sycl::queue &Queue, sycl::kernel &Kernel) { +void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { constexpr int Range = 10; int *usmPtr = sycl::malloc_shared(Range, Queue); int start = 3; @@ -74,15 +74,13 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel) { memset(usmPtr, 0, Range * sizeof(int)); Queue.submit([&](sycl::handler &Handler) { Handler.set_arg(0, usmPtr); - // Handler.set_arg(1, start); - // Handler.set_arg(2, Range); Handler.parallel_for(R1, Kernel); }); Queue.wait(); for (int i = 0; i < Range; i++) { std::cout << usmPtr[i] << " "; - // assert(usmPtr[i] = i + 42); + assert(usmPtr[i] = i + seed); } std::cout << std::endl; @@ -133,22 +131,17 @@ void test_build_and_run() { assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); - // extern "C" was used, so the name "ff_cp" is not mangled + // extern "C" was used, so the name "ff_cp" is not mangled and can be used + // directly. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); - // the templated function name will have been mangled. Mapping from original + // The templated function name will have been mangled. Mapping from original // name to mangled is not yet supported. sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); - // COMING SOON - // sycl::kernel_bundle kb - // = syclexp::build(kb_src, - // syclexp::properties{syclexp::registered_kernel_names{"mykernels::bar"}}); - // sycl::kernel k = kb.ext_oneapi_get_kernel("mykernels::bar"); - - // 4 - test_1(q, k); - test_1(q, k2); + // test the kernels + test_1(q, k, 37 + 5); // AddEm will add 5 more + test_1(q, k2, 39); } void test_error() { From 7eb8f44d9762ce5261717562a48c38380ffc0d39 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 28 Jun 2024 15:31:05 -0700 Subject: [PATCH 25/27] more reviewer feedback --- sycl/include/sycl/kernel_bundle.hpp | 2 + sycl/source/detail/kernel_bundle_impl.hpp | 2 +- .../kernel_compiler/kernel_compiler_sycl.cpp | 246 +++++++++--------- .../KernelCompiler/kernel_compiler_sycl.cpp | 32 ++- 4 files changed, 148 insertions(+), 134 deletions(-) diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp index 18ab1bfe5f994..99663183ad8d3 100644 --- a/sycl/include/sycl/kernel_bundle.hpp +++ b/sycl/include/sycl/kernel_bundle.hpp @@ -874,9 +874,11 @@ struct is_property_key_of struct registered_kernel_names : detail::run_time_property_key { std::vector kernel_names; + registered_kernel_names() {} registered_kernel_names(const std::string &knArg) : kernel_names{knArg} {} registered_kernel_names(const std::vector &knsArg) : kernel_names(knsArg) {} + void add(const std::string &name) { kernel_names.push_back(name); } }; using registered_kernel_names_key = registered_kernel_names; diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index e7052b8b7c405..35a5b690ccb8a 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -739,7 +739,7 @@ class kernel_bundle_impl { bundle_state MState; // ext_oneapi_kernel_compiler : Source, Languauge, KernelNames, IncludePairs - // Language is for both state::source and state::executable + // Language is for both state::source and state::executable. syclex::source_language Language = syclex::source_language::opencl; const std::variant> Source; // only kernel_bundles created from source have KernelNames member. diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 0ca0b313158ca..98f296bb6ceb1 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -29,7 +29,7 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames) { throw sycl::exception(sycl::errc::build, - "kernel_compiler does not supprot GCC<8"); + "kernel_compiler does not support GCC<8"); } } // namespace detail } // namespace ext::oneapi::experimental @@ -50,89 +50,90 @@ inline namespace _V1 { namespace ext::oneapi::experimental { namespace detail { -std::string generate_semi_unique_id() { - // Get the current time as a time_t object - std::time_t now = std::time(nullptr); +std::string generateSemiUniqueId() { + // Get the current time as a time_t object. + std::time_t CurrentTime = std::time(nullptr); - // Convert time_t to a string with format YYYYMMDD_HHMMSS - std::tm *local_time = std::localtime(&now); - std::stringstream ss; - ss << std::put_time(local_time, "%Y%m%d_%H%M%S"); + // Convert time_t to a string with format YYYYMMDD_HHMMSS. + std::tm *LocalTime = std::localtime(&CurrentTime); + std::stringstream Ss; + Ss << std::put_time(LocalTime, "%Y%m%d_%H%M%S"); - // amend with random number - std::random_device rd; - int random_number = rd() % 900 + 100; - ss << "_" << std::setfill('0') << std::setw(3) - << random_number; // Pad with leading zeros + // Amend with random number. + std::random_device Rd; + int RandomNumber = Rd() % 900 + 100; + Ss << "_" << std::setfill('0') << std::setw(3) << RandomNumber; - return ss.str(); + return Ss.str(); } -std::filesystem::path prepare_ws(const std::string &id) { - const std::filesystem::path tmp = std::filesystem::temp_directory_path(); - - std::filesystem::path new_directory_path = tmp / id; +std::filesystem::path prepareWS(const std::string &Id) { + const std::filesystem::path TmpDirectoryPath = + std::filesystem::temp_directory_path(); + std::filesystem::path NewDirectoryPath = TmpDirectoryPath / Id; try { - std::filesystem::create_directories(new_directory_path); - } catch (std::filesystem::filesystem_error const &e) { - throw sycl::exception(sycl::errc::build, e.what()); + std::filesystem::create_directories(NewDirectoryPath); + } catch (const std::filesystem::filesystem_error &E) { + throw sycl::exception(sycl::errc::build, E.what()); } - return new_directory_path; + return NewDirectoryPath; } -std::string user_args_as_string(const std::vector &UserArgs) { - return std::accumulate(UserArgs.begin(), UserArgs.end(), std::string(""), - [](const std::string &a, const std::string &b) { - return a.empty() ? b : a + " " + b; +std::string userArgsAsString(const std::vector &UserArguments) { + return std::accumulate(UserArguments.begin(), UserArguments.end(), + std::string(""), + [](const std::string &A, const std::string &B) { + return A.empty() ? B : A + " " + B; }); } -void output_preamble(std::ofstream &os, const std::filesystem::path &file_path, - const std::string &id, - const std::vector &UserArgs) { +void outputPreamble(std::ofstream &Os, const std::filesystem::path &FilePath, + const std::string &Id, + const std::vector &UserArgs) { - os << "/*\n clang++ -fsycl -o " << id << ".bin " - << user_args_as_string(UserArgs) - << "-fno-sycl-dead-args-optimization -fsycl-dump-device-code=./ " << id - << ".cpp \n */" << std::endl; + Os << "/*\n"; + Os << " clang++ -fsycl -o " << Id << ".bin "; + Os << userArgsAsString(UserArgs); + Os << " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=./ " << Id; + Os << ".cpp \n */" << std::endl; } std::filesystem::path -output_cpp(const std::filesystem::path &parent_dir, const std::string &id, - std::string raw_code_string, - const std::vector &UserArgs, - const std::vector &RegisteredKernelNames) { - std::filesystem::path file_path = parent_dir / (id + ".cpp"); - std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); - - if (outfile.is_open()) { - output_preamble(outfile, file_path, id, UserArgs); - outfile << raw_code_string << std::endl; - - // temporarily needed until -c works with -fsycl-dump-spirv - outfile << "int main(){\n"; - for (std::string nm : RegisteredKernelNames) { - outfile << " " << nm << ";\n"; +outputCpp(const std::filesystem::path &ParentDir, const std::string &Id, + std::string RawCodeString, const std::vector &UserArgs, + const std::vector &RegisteredKernelNames) { + std::filesystem::path FilePath = ParentDir / (Id + ".cpp"); + std::ofstream Outfile(FilePath, std::ios::out | std::ios::trunc); + + if (Outfile.is_open()) { + outputPreamble(Outfile, FilePath, Id, UserArgs); + Outfile << RawCodeString << std::endl; + + // Temporarily needed until -c works with -fsycl-dump-spirv. + Outfile << "int main() {\n"; + for (const std::string &KernelName : RegisteredKernelNames) { + Outfile << " " << KernelName << ";\n"; } - outfile << " return 0;\n}" << std::endl; + Outfile << " return 0;\n}\n" << std::endl; - outfile.close(); // Close the file when finished + Outfile.close(); } else { throw sycl::exception(sycl::errc::build, "Failed to open .cpp file for write: " + - file_path.string()); + FilePath.string()); } - return file_path; + + return FilePath; } -void output_include_files(const std::filesystem::path &dpath, - include_pairs_t IncludePairs) { +void outputIncludeFiles(const std::filesystem::path &Dirpath, + include_pairs_t IncludePairs) { using pairStrings = std::pair; for (pairStrings p : IncludePairs) { - std::filesystem::path file_path = dpath / p.first; - std::ofstream outfile(file_path, std::ios::out | std::ios::trunc); + std::filesystem::path FilePath = Dirpath / p.first; + std::ofstream outfile(FilePath, std::ios::out | std::ios::trunc); if (outfile.is_open()) { outfile << p.second << std::endl; @@ -140,92 +141,93 @@ void output_include_files(const std::filesystem::path &dpath, } else { throw sycl::exception(sycl::errc::build, "Failed to open include file for write: " + - file_path.string()); + FilePath.string()); } } } -std::string get_compiler_name() { +std::string getCompilerName() { #ifdef __WIN32 - std::string compiler = "clang++.exe"; + std::string Compiler = "clang++.exe"; #else - std::string compiler = "clang++"; + std::string Compiler = "clang++"; #endif - return compiler; + return Compiler; } -void invoke_compiler(const std::filesystem::path &fpath, - const std::filesystem::path &dpath, const std::string &id, - const std::vector &UserArgs, - std::string *LogPtr) { +void invokeCompiler(const std::filesystem::path &FPath, + const std::filesystem::path &DPath, const std::string &Id, + const std::vector &UserArgs, + std::string *LogPtr) { - std::filesystem::path file_path(fpath); - std::filesystem::path parent_dir(dpath); - std::filesystem::path target_path = parent_dir / (id + ".bin"); - std::filesystem::path log_path = parent_dir / "compilation_log.txt"; - std::string compiler = get_compiler_name(); + std::filesystem::path FilePath(FPath); + std::filesystem::path ParentDir(DPath); + std::filesystem::path TargetPath = ParentDir / (Id + ".bin"); + std::filesystem::path LogPath = ParentDir / "compilation_log.txt"; + std::string Compiler = getCompilerName(); - std::string command = - compiler + " -fsycl -o " + target_path.make_preferred().string() + " " + - user_args_as_string(UserArgs) + + std::string Command = + Compiler + " -fsycl -o " + TargetPath.make_preferred().string() + " " + + userArgsAsString(UserArgs) + " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=" + - parent_dir.make_preferred().string() + " " + - file_path.make_preferred().string() + " 2> " + - log_path.make_preferred().string(); + ParentDir.make_preferred().string() + " " + + FilePath.make_preferred().string() + " 2> " + + LogPath.make_preferred().string(); - int result = std::system(command.c_str()); + int Result = std::system(Command.c_str()); - // Read the log file contents into the log variable + // Read the log file contents into the log variable. std::string CompileLog; - std::ifstream log_stream; - log_stream.open(log_path); - if (log_stream.is_open()) { - std::stringstream log_buffer; - log_buffer << log_stream.rdbuf(); - CompileLog.append(log_buffer.str()); + std::ifstream LogStream; + LogStream.open(LogPath); + if (LogStream.is_open()) { + std::stringstream LogBuffer; + LogBuffer << LogStream.rdbuf(); + CompileLog.append(LogBuffer.str()); if (LogPtr != nullptr) - LogPtr->append(log_buffer.str()); + LogPtr->append(LogBuffer.str()); - } else if (result == 0 && LogPtr != nullptr) { - // if there was a compilation problem, we want to report that (below) + } else if (Result == 0 && LogPtr != nullptr) { + // If there was a compilation problem, we want to report that (below), // not a mere "missing log" error. throw sycl::exception(sycl::errc::build, "failure retrieving compilation log"); } - if (result != 0) { + if (Result != 0) { throw sycl::exception(sycl::errc::build, - "Compile failure: " + std::to_string(result) + " " + + "Compile failure: " + std::to_string(Result) + " " + CompileLog); } } -std::filesystem::path find_spv(const std::filesystem::path &parent_dir, - const std::string &id) { - std::regex pattern_regex(id + R"(.*\.spv)"); +std::filesystem::path findSpv(const std::filesystem::path &ParentDir, + const std::string &Id) { + std::regex PatternRegex(Id + R"(.*\.spv)"); - // Iterate through all files in the directory matching the pattern - for (const auto &entry : std::filesystem::directory_iterator(parent_dir)) { - if (entry.is_regular_file() && - std::regex_match(entry.path().filename().string(), pattern_regex)) { - return entry.path(); + // Iterate through all files in the directory matching the pattern. + for (const auto &Entry : std::filesystem::directory_iterator(ParentDir)) { + if (Entry.is_regular_file() && + std::regex_match(Entry.path().filename().string(), PatternRegex)) { + return Entry.path(); // Return the path if it matches the SPV pattern. } } - // File not found, throw - throw sycl::exception(sycl::errc::build, "SPIRV output matching " + id + + + // File not found, throw. + throw sycl::exception(sycl::errc::build, "SPIRV output matching " + Id + " missing from " + - parent_dir.filename().string()); + ParentDir.filename().string()); } -spirv_vec_t load_spv_from_file(std::filesystem::path file_name) { - std::ifstream spv_stream(file_name, std::ios::binary); - spv_stream.seekg(0, std::ios::end); - size_t sz = spv_stream.tellg(); - spv_stream.seekg(0); - spirv_vec_t spv(sz); - spv_stream.read(reinterpret_cast(spv.data()), sz); +spirv_vec_t loadSpvFromFile(const std::filesystem::path &FileName) { + std::ifstream SpvStream(FileName, std::ios::binary); + SpvStream.seekg(0, std::ios::end); + size_t Size = SpvStream.tellg(); + SpvStream.seekg(0); + spirv_vec_t Spv(Size); + SpvStream.read(reinterpret_cast(Spv.data()), Size); - return spv; + return Spv; } spirv_vec_t @@ -233,25 +235,25 @@ SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames) { // clang-format off - const std::string id = generate_semi_unique_id(); - const std::filesystem::path parent_dir = prepare_ws(id); - std::filesystem::path file_path = output_cpp(parent_dir, id, SYCLSource, UserArgs, RegisteredKernelNames); - output_include_files(parent_dir, IncludePairs); - invoke_compiler(file_path, parent_dir, id, UserArgs, LogPtr); - std::filesystem::path spv_path = find_spv(parent_dir, id); - return load_spv_from_file(spv_path); + const std::string id = generateSemiUniqueId(); + const std::filesystem::path ParentDir = prepareWS(id); + std::filesystem::path FilePath = outputCpp(ParentDir, id, SYCLSource, UserArgs, RegisteredKernelNames); + outputIncludeFiles(ParentDir, IncludePairs); + invokeCompiler(FilePath, ParentDir, id, UserArgs, LogPtr); + std::filesystem::path SpvPath = findSpv(ParentDir, id); + return loadSpvFromFile(SpvPath); // clang-format on } bool SYCL_Compilation_Available() { - // is compiler on $PATH ? We try to invoke it. - std::string id = generate_semi_unique_id(); + // Is compiler on $PATH ? We try to invoke it. + std::string id = generateSemiUniqueId(); const std::filesystem::path tmp = std::filesystem::temp_directory_path(); - std::filesystem::path dump_path = tmp / (id + "_version.txt"); - std::string compiler = get_compiler_name(); - std::string test_command = - compiler + " --version &> " + dump_path.make_preferred().string(); - int result = std::system(test_command.c_str()); + std::filesystem::path DumpPath = tmp / (id + "_version.txt"); + std::string Compiler = getCompilerName(); + std::string TestCommand = + Compiler + " --version &> " + DumpPath.make_preferred().string(); + int result = std::system(TestCommand.c_str()); return (result == 0); } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index ebf0e15ce1751..b140deb8769b3 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -105,21 +105,21 @@ void test_build_and_run() { return; } - // create from source + // Create from source. source_kb kbSrc = syclex::create_kernel_bundle_from_source( ctx, syclex::source_language::sycl, SYCLSource, syclex::properties{syclex::include_files{"AddEm.h", AddEmH}}); - // double check kernel_bundle.get_source() / get_backend() + // Double check kernel_bundle.get_source() / get_backend(). sycl::context ctxRes = kbSrc.get_context(); assert(ctxRes == ctx); sycl::backend beRes = kbSrc.get_backend(); assert(beRes == ctx.get_backend()); - // compilation of empty prop list, no devices + // Compilation of empty prop list, no devices. exe_kb kbExe1 = syclex::build(kbSrc); - // compilation with props and devices + // Compilation with props and devices std::string log; std::vector flags{"-g", "-fno-fast-math"}; std::vector devs = kbSrc.get_devices(); @@ -131,16 +131,23 @@ void test_build_and_run() { assert(log.find("warning: 'this_nd_item<1>' is deprecated") != std::string::npos); - // extern "C" was used, so the name "ff_cp" is not mangled and can be used - // directly. + // clang-format off + + // extern "C" was used, so the name "ff_cp" is not mangled and can be used directly. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); + // The templated function name will have been mangled. Mapping from original - // name to mangled is not yet supported. - sycl::kernel k2 = - kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); + // name to mangled is not yet supported. So we cannot yet do this: + // sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); + + // Instead, we can TEMPORARILY use the mangled name. Once demangling is supported + // this might no longer work. + sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); + + // clang-format on - // test the kernels - test_1(q, k, 37 + 5); // AddEm will add 5 more + // Test the kernels. + test_1(q, k, 37 + 5); // AddEm will add 5 more. test_1(q, k2, 39); } @@ -166,6 +173,9 @@ void test_error() { } catch (sycl::exception &e) { // yas! assert(e.code() == sycl::errc::build); + assert(std::string(e.what()).find( + "error: expected ';' at end of declaration") != + std::string::npos); } } From 132525c0b6b24778f484479100ded85194a47a2e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 1 Jul 2024 14:24:11 -0700 Subject: [PATCH 26/27] add esimd_kernel to testing --- .../KernelCompiler/kernel_compiler_sycl.cpp | 85 +++++++++++++++++++ 1 file changed, 85 insertions(+) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index b140deb8769b3..fc294f49fad45 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -62,6 +62,27 @@ void ff_cp(int *ptr) { } )==="; +auto constexpr ESIMDSource = R"===( +#include +#include + +using namespace sycl::ext::intel::esimd; + +constexpr int VL = 16; + +extern "C" SYCL_EXTERNAL SYCL_ESIMD_KERNEL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void vector_add_esimd(float *A, float *B, float *C) { + sycl::nd_item<1> item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + unsigned int i = item.get_global_id(0); + unsigned int offset = i * VL ; + + simd va(A + offset); + simd vb(B + offset); + simd vc = va + vb; + vc.copy_to(C + offset); +} +)==="; + void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { constexpr int Range = 10; int *usmPtr = sycl::malloc_shared(Range, Queue); @@ -179,11 +200,75 @@ void test_error() { } } +void test_esimd() { + 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(); + + if (!q.get_device().has(sycl::aspect::ext_intel_esimd)) { + std::cout << "Device does not support ESIMD" << std::endl; + return; + } + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return; + } + + std::string log; + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, ESIMDSource); + exe_kb kbExe = + syclex::build(kbSrc, syclex::properties{syclex::save_log{&log}}); + + // extern "C" was used, so the name "vector_add_esimd" is not mangled and can + // be used directly. + sycl::kernel k = kbExe.ext_oneapi_get_kernel("vector_add_esimd"); + + // Now test it. + constexpr int VL = 16; // this constant also in ESIMDSource string. + constexpr int size = VL * 16; + + float *A = sycl::malloc_shared(size, q); + float *B = sycl::malloc_shared(size, q); + float *C = sycl::malloc_shared(size, q); + for (size_t i = 0; i < size; i++) { + A[i] = float(1); + B[i] = float(2); + C[i] = 0.0f; + } + sycl::range<1> GlobalRange{size / VL}; + sycl::range<1> LocalRange{1}; + sycl::nd_range<1> NDRange{GlobalRange, LocalRange}; + + q.submit([&](sycl::handler &h) { + h.set_arg(0, A); + h.set_arg(1, B); + h.set_arg(2, C); + h.parallel_for(NDRange, k); + }).wait(); + + // Check. + for (size_t i = 0; i < size; i++) { + assert(C[i] == 3.0f); + } + + sycl::free(A, q); + sycl::free(B, q); + sycl::free(C, q); +} + int main() { #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER test_build_and_run(); test_error(); + test_esimd(); #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif From b71925c545f69e9b6f5b0603df70b99f5cec9059 Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Mon, 1 Jul 2024 17:05:49 -0700 Subject: [PATCH 27/27] win symbols --- sycl/test/abi/sycl_symbols_windows.dump | 31 ++++++++++--------------- 1 file changed, 12 insertions(+), 19 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 37a3743a81255..00d45e8913778 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -569,10 +569,10 @@ ??0half@host_half_impl@detail@_V1@sycl@@QEAA@AEBM@Z ??0half@host_half_impl@detail@_V1@sycl@@QEAA@G@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z -??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N1@Z -??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z +??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@00_N@Z ??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N1@Z +??0handler@_V1@sycl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_N@Z ??0host_selector@_V1@sycl@@QEAA@$$QEAV012@@Z ??0host_selector@_V1@sycl@@QEAA@AEBV012@@Z ??0host_selector@_V1@sycl@@QEAA@XZ @@ -609,10 +609,6 @@ ??0kernel_id@_V1@sycl@@AEAA@PEBD@Z ??0kernel_id@_V1@sycl@@QEAA@$$QEAV012@@Z ??0kernel_id@_V1@sycl@@QEAA@AEBV012@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVqueue@45@_K@Z -??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA@AEBV?$shared_ptr@Vgraph_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z @@ -621,6 +617,10 @@ ??0node@experimental@oneapi@ext@_V1@sycl@@AEAA@AEBV?$shared_ptr@Vnode_impl@detail@experimental@oneapi@ext@_V1@sycl@@@std@@@Z ??0node@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z ??0node@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVdevice@45@AEBVcontext@45@_K@Z +??0physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBVqueue@45@_K@Z ??0platform@_V1@sycl@@AEAA@AEBVdevice@12@@Z ??0platform@_V1@sycl@@AEAA@V?$shared_ptr@Vplatform_impl@detail@_V1@sycl@@@std@@@Z ??0platform@_V1@sycl@@QEAA@$$QEAV012@@Z @@ -683,9 +683,9 @@ ??1kernel@_V1@sycl@@QEAA@XZ ??1kernel_bundle_plain@detail@_V1@sycl@@QEAA@XZ ??1kernel_id@_V1@sycl@@QEAA@XZ -??1physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1node@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ +??1physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1platform@_V1@sycl@@QEAA@XZ ??1queue@_V1@sycl@@QEAA@XZ ??1sampler@_V1@sycl@@QEAA@XZ @@ -768,12 +768,12 @@ ??4kernel_bundle_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4kernel_id@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4kernel_id@_V1@sycl@@QEAAAEAV012@AEBV012@@Z -??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z -??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@$$QEAV0123456@@Z ??4modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV0123456@AEBV0123456@@Z ??4node@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ??4node@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z +??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z +??4physical_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4platform@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4platform@_V1@sycl@@QEAAAEAV012@AEBV012@@Z ??4queue@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z @@ -3985,10 +3985,6 @@ ?ext_intel_read_host_pipe@handler@_V1@sycl@@AEAAXVstring_view@detail@23@PEAX_K_N@Z ?ext_intel_write_host_pipe@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEAX_K_N@Z ?ext_intel_write_host_pipe@handler@_V1@sycl@@AEAAXVstring_view@detail@23@PEAX_K_N@Z -<<<<<<< HEAD -?ext_oneapi_advise_usm_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEBX_KW4_pi_mem_advice@@V?$vector@IV?$allocator@I@std@@@6@PEAI@Z -======= ->>>>>>> sycl ?ext_oneapi_architecture_is@device@_V1@sycl@@QEAA_NW4arch_category@experimental@oneapi@ext@23@@Z ?ext_oneapi_architecture_is@device@_V1@sycl@@QEAA_NW4architecture@experimental@oneapi@ext@23@@Z ?ext_oneapi_barrier@handler@_V1@sycl@@QEAAXAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@@Z @@ -4059,7 +4055,6 @@ ?ext_oneapi_owner_before@?$OwnerLessBase@Vqueue@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVqueue@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vstream@_V1@sycl@@@2oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVstream@34@@Z -?ext_oneapi_prefetch_usm_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAX_KV?$vector@IV?$allocator@I@std@@@6@PEAI@Z ?ext_oneapi_prod@queue@_V1@sycl@@QEAAXXZ ?ext_oneapi_set_external_event@queue@_V1@sycl@@QEAAXAEBVevent@23@@Z ?ext_oneapi_signal_external_semaphore@handler@_V1@sycl@@QEAAXUinterop_semaphore_handle@experimental@oneapi@ext@23@@Z @@ -4091,14 +4086,13 @@ ?find_device_intersection@detail@_V1@sycl@@YA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV?$vector@V?$kernel_bundle@$00@_V1@sycl@@V?$allocator@V?$kernel_bundle@$00@_V1@sycl@@@std@@@5@@Z ?free@_V1@sycl@@YAXPEAXAEBVcontext@12@AEBUcode_location@detail@12@@Z ?free@_V1@sycl@@YAXPEAXAEBVqueue@12@AEBUcode_location@detail@12@@Z -?free_virtual_mem@experimental@oneapi@ext@_V1@sycl@@YAX_K0AEBVcontext@45@@Z +?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@W4image_type@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@W4image_type@12345@AEBVqueue@45@@Z -?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z ?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z -?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?free_virtual_mem@experimental@oneapi@ext@_V1@sycl@@YAX_K0AEBVcontext@45@@Z ?frexp_impl@detail@_V1@sycl@@YA?AVhalf@half_impl@123@V45123@PEAH@Z ?frexp_impl@detail@_V1@sycl@@YAMMPEAH@Z ?frexp_impl@detail@_V1@sycl@@YANNPEAH@Z @@ -4213,8 +4207,8 @@ ?get_coordinate_normalization_mode@sampler@_V1@sycl@@QEBA?AW4coordinate_normalization_mode@23@XZ ?get_count@image_plain@detail@_V1@sycl@@IEBA_KXZ ?get_descriptor@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBAAEBUimage_descriptor@23456@XZ -?get_device@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ ?get_device@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ +?get_device@physical_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ ?get_device@queue@_V1@sycl@@QEBA?AVdevice@23@XZ ?get_devices@context@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?get_devices@device@_V1@sycl@@SA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@W4device_type@info@23@@Z @@ -4482,7 +4476,6 @@ ?updateAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBVAccessorBaseHost@267@@Z ?updateValue@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBX_K@Z ?use_kernel_bundle@handler@_V1@sycl@@QEAAXAEBV?$kernel_bundle@$01@23@@Z -?usesPinnedHostMemory@SYCLMemObjT@detail@_V1@sycl@@UEBA_NXZ ?verifyDeviceHasProgressGuarantee@handler@_V1@sycl@@AEAAXW4forward_progress_guarantee@experimental@oneapi@ext@23@W4execution_scope@56723@1@Z ?verifyKernelInvoc@handler@_V1@sycl@@AEAAXAEBVkernel@23@@Z ?verifyUsedKernelBundle@handler@_V1@sycl@@AEAAXAEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z