From b1c348e1b96ca21ad11ef723413f72e2b1eca27a Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Mon, 30 Nov 2020 17:23:40 -0800 Subject: [PATCH 01/10] [SYCL] Implement INTEL feature class online_compiler Signed-off-by: Vyacheslav N Klochkov --- .../include/CL/sycl/INTEL/online_compiler.hpp | 119 ++++++---- sycl/include/CL/sycl/exception.hpp | 2 +- sycl/source/CMakeLists.txt | 1 + .../source/detail/online_compiler/ocloc_api.h | 110 +++++++++ .../online_compiler/online_compiler.cpp | 208 ++++++++++++++++++ sycl/test/abi/sycl_symbols_linux.dump | 4 + .../online_compiler/online_compiler_L0.cpp | 46 ++++ .../online_compiler_OpenCL.cpp | 33 +++ .../online_compiler_common.hpp | 171 ++++++++++++++ 9 files changed, 648 insertions(+), 46 deletions(-) create mode 100644 sycl/source/detail/online_compiler/ocloc_api.h create mode 100644 sycl/source/detail/online_compiler/online_compiler.cpp create mode 100644 sycl/test/on-device/online_compiler/online_compiler_L0.cpp create mode 100644 sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp create mode 100644 sycl/test/on-device/online_compiler/online_compiler_common.hpp diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index a2e8c6ba5b41..bbef18c76dda 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -8,11 +8,11 @@ #pragma once -#include #include // for __SYCL_INLINE_NAMESPACE +#include // for __SYCL_EXPORT #include -#include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -59,7 +59,14 @@ class device_arch { /// Represents an error happend during online compilation. class online_compile_error : public sycl::exception { - // TBD +public: + online_compile_error() = default; + online_compile_error(const string_class &Msg) : Msg(Msg) {} + + const char *what() const noexcept override { return Msg.c_str(); }; + +private: + string_class Msg; }; /// Designates a source language for the online compiler. @@ -67,28 +74,35 @@ enum class source_language { opencl_c, cm }; /// Represents an online compiler for the language given as template /// parameter. -template class online_compiler { +template class __SYCL_EXPORT online_compiler { public: /// Constructs online compiler which can target any device and produces - /// given compiled code format. Produces device code is 64-bit. + /// given compiled code format. Produces 64-bit device code. /// The created compiler is "optimistic" - it assumes all applicable SYCL /// device capabilities are supported by the target device(s). online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) : OutputFormat(fmt), OutputFormatVersion({0, 0}), + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler which targets given architecture and produces - /// given compiled code format. Produces device code is 64-bit. + /// given compiled code format. Produces 64-bit device code. /// Throws online_compile_error if values of constructor arguments are /// contradictory or not supported - e.g. if the source language is not /// supported for given device type. online_compiler(sycl::info::device_type dev_type, device_arch arch, compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceArch(arch), - Is64Bit(true), DeviceStepping("") {} + : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), + DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler for the target specified by given SYCL device. - online_compiler(const sycl::device &dev); + // TODO: the initial version generates the generic code (SKL now), need + // to do additional device::info calls to determine the device by it's + // features. + online_compiler(const sycl::device &dev) + : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} /// Compiles given in-memory \c Lang source to a binary blob. Blob format, /// other parameters are set in the constructor by the compilation target @@ -100,31 +114,50 @@ template class online_compiler { std::vector compile(const std::string &src, const Tys &... args); /// Sets the compiled code format of the compilation target and returns *this. - online_compiler &setOutputFormat(compiled_code_format fmt); + online_compiler &setOutputFormat(compiled_code_format fmt) { + OutputFormat = fmt; + return *this; + } /// Sets the compiled code format version of the compilation target and /// returns *this. - online_compiler &setOutputFormatVersion(int major, int minor); + online_compiler &setOutputFormatVersion(int major, int minor) { + OutputFormatVersion = {major, minor}; + return *this; + } /// Sets the device type of the compilation target and returns *this. - online_compiler &setTargetDeviceType(sycl::info::device_type type); + online_compiler &setTargetDeviceType(sycl::info::device_type type) { + DeviceType = type; + return *this; + } /// Sets the device architecture of the compilation target and returns *this. - online_compiler &setTargetDeviceArch(device_arch arch); + online_compiler &setTargetDeviceArch(device_arch arch) { + DeviceArch = arch; + return *this; + } /// Makes the compilation target 32-bit and returns *this. - online_compiler &set32bitTarget(); + online_compiler &set32bitTarget() { + Is64Bit = false; + return *this; + }; /// Makes the compilation target 64-bit and returns *this. - online_compiler &set64bitTarget(); + online_compiler &set64bitTarget() { + Is64Bit = true; + return *this; + }; /// Sets implementation-defined target device stepping of the compilation /// target and returns *this. - online_compiler &setTargetDeviceStepping(const std::string &id); + online_compiler &setTargetDeviceStepping(const std::string &id) { + DeviceStepping = id; + return *this; + } private: - // Compilation target specification fields: { - /// Compiled code format. compiled_code_format OutputFormat; @@ -142,51 +175,47 @@ template class online_compiler { /// Target device stepping (implementation defined) std::string DeviceStepping; - // } + + /// Handles to helper functions used by the implementation. + void *CompileToSPIRVHandle = nullptr; + void *FreeSPIRVOutputsHandle = nullptr; }; // Specializations of the online_compiler class and 'compile' function for // particular languages and parameter types. -/// Compiles given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source -template <> -template <> -std::vector -online_compiler::compile(const std::string &src) { - // real implementation will call some non-templated impl function here - return std::vector{}; -} - -/// Compiles given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. /// @param options - compilation options (implementation defined); standard -/// OpenCL JIT compiler options must be supported +/// OpenCL JIT compiler options must be supported. template <> template <> std::vector online_compiler::compile( - const std::string &src, const std::vector &options) { - // real implementation will call some non-templated impl function here - return std::vector{}; -} + const std::string &src, const std::vector &options); -/// Compiles given CM source. +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. template <> template <> std::vector -online_compiler::compile(const std::string &src) { - // real implementation will call some non-templated impl function here - return std::vector{}; +online_compiler::compile(const std::string &src) { + return compile(src, std::vector{}); } -/// Compiles given CM source. -/// @param options - compilation options (implementation defined) +/// Compiles the given CM source \p src. +/// @param src - contents of the source. +/// @param options - compilation options (implementation defined). template <> template <> std::vector online_compiler::compile( - const std::string &src, const std::vector &options) { - // real implementation will call some non-templated impl function here - return std::vector{}; + const std::string &src, const std::vector &options); + +/// Compiles the given CM source \p src. +template <> +template <> +std::vector +online_compiler::compile(const std::string &src) { + return compile(src, std::vector{}); } } // namespace INTEL diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp index 0004bc311a4c..90def5a10301 100644 --- a/sycl/include/CL/sycl/exception.hpp +++ b/sycl/include/CL/sycl/exception.hpp @@ -30,7 +30,7 @@ class __SYCL_EXPORT exception : public std::exception { public: exception() = default; - const char *what() const noexcept final override; + const char *what() const noexcept override; bool has_context() const; diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index b1eeba9f8b2e..b02fcb00b74a 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -132,6 +132,7 @@ set(SYCL_SOURCES "detail/program_impl.cpp" "detail/program_manager/program_manager.cpp" "detail/queue_impl.cpp" + "detail/online_compiler/online_compiler.cpp" "detail/os_util.cpp" "detail/platform_util.cpp" "detail/reduction.cpp" diff --git a/sycl/source/detail/online_compiler/ocloc_api.h b/sycl/source/detail/online_compiler/ocloc_api.h new file mode 100644 index 000000000000..f073c5c524f7 --- /dev/null +++ b/sycl/source/detail/online_compiler/ocloc_api.h @@ -0,0 +1,110 @@ +//===------- ocloc_api.h --------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// This file is copied from +// https://github.com/intel/compute-runtime/blob/master/shared/offline_compiler/source/ocloc_api.h + +#include + +#ifndef OCLOC_MAKE_VERSION +/// Generates ocloc API versions +#define OCLOC_MAKE_VERSION(_major, _minor) \ + ((_major << 16) | (_minor & 0x0000ffff)) +#endif // OCLOC_MAKE_VERSION + +typedef enum _ocloc_version_t { + OCLOC_VERSION_1_0 = OCLOC_MAKE_VERSION(1, 0), ///< version 1.0 + OCLOC_VERSION_CURRENT = OCLOC_MAKE_VERSION(1, 0), ///< latest known version + OCLOC_VERSION_FORCE_UINT32 = 0x7fffffff +} ocloc_version_t; + +#ifdef _WIN32 +#define SIGNATURE __declspec(dllexport) int __cdecl +#else +#define SIGNATURE int +#endif + +extern "C" { +/// Invokes ocloc API using C interface. Supported commands match +/// the functionality of ocloc executable (check ocloc's "help" +/// for reference : shared/offline_compiler/source/ocloc_api.cpp) +/// at https://github.com/intel/compute-runtime. +/// +/// NumArgs and argv params represent the command line. +/// Remaining params represent I/O. +/// Output params should be freed using oclocFreeOutput() when +/// no longer needed. +/// List and names of outputs match outputs of ocloc executable. +/// +/// \param NumArgs is the number of arguments to pass to ocloc. +/// +/// \param Argv is an array of arguments to be passed to ocloc. +/// +/// \param NumSources is the number of in-memory representations +/// of source files to be passed to ocloc. +/// +/// \param DataSources is an array of in-memory representations +/// of source files to be passed to ocloc. +/// +/// \param LenSources is an array of sizes of in-memory representations +/// of source files passed to ocloc as DataSources. +/// +/// \param NameSources is an array of names of in-memory representations +/// of source files passed to ocloc as DataSources. +/// +/// \param NumInputHeaders is the number of in-memory representations +/// of header files to be passed to ocloc. +/// +/// \param DataInputHeaders is an array of in-memory representations +/// of header files to be passed to ocloc. +/// +/// \param LenInputHeaders is an array of sizes of in-memory representations +/// of header files passed to ocloc as DataInputHeaders. +/// +/// \param NameInputHeaders is an array of names of in-memory representations +/// of header files passed to ocloc as DataInputHeaders. +/// +/// \param NumOutputs returns the number of outputs. +/// +/// \param DataOutputs returns an array of in-memory representations +/// of output files. +/// +/// \param LenOutputs returns an array of sizes of in-memory representations +/// of output files. +/// +/// \param NameOutputs returns an array of names of in-memory representations +/// of output files. Special name stdout.log describes output that contains +/// messages generated by ocloc (e.g. compiler errors/warnings). +/// +/// \returns 0 on succes. Returns non-0 in case of failure. +SIGNATURE oclocInvoke(uint32_t NumArgs, const char *Argv[], uint32_t NumSources, + const uint8_t **DataSources, const uint64_t *LenSources, + const char **NameSources, uint32_t NumInputHeaders, + const uint8_t **DataInputHeaders, + const uint64_t *LenInputHeaders, + const char **NameInputHeaders, uint32_t *NumOutputs, + uint8_t ***DataOutputs, uint64_t **LenOutputs, + char ***NameOutputs); + +/// Frees results of oclocInvoke +/// +/// \param NumOutputs is number of outputs as returned by oclocInvoke(). +/// +/// \param DataOutputs is array of outputs as returned by oclocInvoke(). +/// +/// \param LenOutputs is array of sizes of outputs as returned by oclocInvoke(). +/// +/// \param NameOutputs is array of names of outputs as returned by oclocInvoke() +/// +/// \returns 0 on succes. Returns non-0 in case of failure. +SIGNATURE oclocFreeOutput(uint32_t *NumOutputs, uint8_t ***DataOutputs, + uint64_t **LenOutputs, char ***NameOutputs); + +/// Returns the current version of ocloc. +SIGNATURE oclocVersion(); +} diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp new file mode 100644 index 000000000000..e9187ab469c6 --- /dev/null +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -0,0 +1,208 @@ +//==----------- online_compiler.cpp ----------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include // for __SYCL_INLINE_NAMESPACE +#include +#include + +#include + +#include "ocloc_api.h" + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace INTEL { +namespace detail { + +static std::vector +prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, + bool Is64Bit, const std::string &DeviceStepping, + const std::string &UserArgs) { + std::vector Args = {"ocloc", "-q", "-spv_only", "-device"}; + + if (DeviceType == sycl::info::device_type::gpu) { + switch (DeviceArch) { + case device_arch::gpu_gen9_5: + Args.push_back("cfl"); + break; + + case device_arch::gpu_gen11: + Args.push_back("icllp"); + break; + + default: + Args.push_back("skl"); + } + } else { + Args.push_back("skl"); + } + + if (DeviceStepping != "") { + Args.push_back("-revision_id"); + Args.push_back(DeviceStepping.c_str()); + } + + Args.push_back(Is64Bit ? "-64" : "-32"); + + if (UserArgs != "") { + Args.push_back("-options"); + Args.push_back(UserArgs.c_str()); + } + + return Args; +} + +static std::vector +compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, + const std::string &DeviceStepping, void *&CompileToSPIRVHandle, + void *&FreeSPIRVOutputsHandle, + const std::vector &UserArgs) { + + if (!CompileToSPIRVHandle) { +#ifdef __SYCL_RT_OS_WINDOWS + std::string OclocLibraryName = "ocloc64.dll"; +#else + std::string OclocLibraryName = "libocloc.so"; +#endif + void *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); + if (!OclocLibrary) + throw online_compile_error("Cannot load ocloc library"); + + void *OclocVersionHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion"); + // The initial versions of ocloc library did not have the oclocVersion() + // function. Those versions had the same API as the first version of ocloc + // library having that oclocVersion() function. + int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0; + if (OclocVersionHandle) { + decltype(::oclocVersion) *OclocVersionFunc = + reinterpret_cast(OclocVersionHandle); + LoadedVersion = OclocVersionFunc(); + } + int LoadedVersionMajor = LoadedVersion >> 16; + int CurrentVersionMajor = (ocloc_version_t::OCLOC_VERSION_CURRENT) >> 16; + if (LoadedVersionMajor != CurrentVersionMajor) + throw online_compile_error( + std::string("Found incompatible version of ocloc library: (") + + std::to_string(LoadedVersionMajor) + ", " + + std::to_string(LoadedVersion & 0xffff) + + "). The supported versions are (" + + std::to_string(CurrentVersionMajor) + ", *)."); + + CompileToSPIRVHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); + if (!CompileToSPIRVHandle) + throw online_compile_error("Cannot load oclocInvoke() function"); + FreeSPIRVOutputsHandle = sycl::detail::pi::getOsLibraryFuncAddress( + OclocLibrary, "oclocFreeOutput"); + if (!FreeSPIRVOutputsHandle) + throw online_compile_error("Cannot load oclocFreeOutput() function"); + } + + std::string CombinedUserArgs; + for (auto UserArg : UserArgs) { + if (UserArg == "") + continue; + if (CombinedUserArgs != "") + CombinedUserArgs = CombinedUserArgs + " " + UserArg; + else + CombinedUserArgs = UserArg; + } + std::vector Args = detail::prepareOclocArgs( + DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs); + + uint32_t NumOutputs = 0; + byte **Outputs = nullptr; + size_t *OutputLengths = nullptr; + char **OutputNames = nullptr; + + const byte *Sources[] = {reinterpret_cast(Source.c_str())}; + const char *SourceName = "main.cl"; + const uint64_t SourceLengths[] = {Source.length() + 1}; + + Args.push_back("-file"); + Args.push_back(SourceName); + + decltype(::oclocInvoke) *OclocInvokeFunc = + reinterpret_cast(CompileToSPIRVHandle); + int Error = + OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths, + &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs, + &Outputs, &OutputLengths, &OutputNames); + + byte *SpirV = nullptr; + std::string CompileLog; + size_t SpirVSize = 0; + for (uint32_t I = 0; I < NumOutputs; I++) { + size_t NameLen = strlen(OutputNames[I]); + if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr) { + SpirVSize = OutputLengths[I]; + SpirV = new byte[SpirVSize]; + std::memcpy(SpirV, Outputs[I], SpirVSize); + } else if (!strcmp(OutputNames[I], "stdout.log")) { + CompileLog = std::string(reinterpret_cast(Outputs[I])); + } + } + + if (Error) + throw online_compile_error("ocloc reported compilation errors {" + + CompileLog + "}"); + + decltype(::oclocFreeOutput) *OclocFreeOutputFunc = + reinterpret_cast(FreeSPIRVOutputsHandle); + Error = + OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames); + + if (Error) + throw online_compile_error("ocloc cannot safely free resources"); + + return std::vector(SpirV, SpirV + SpirVSize); +} +} // namespace detail + +template <> +template <> +std::vector online_compiler::compile( + const std::string &Source, const std::vector &UserArgs) { + + if (OutputFormatVersion != std::pair{0, 0}) { + std::string Version = std::to_string(OutputFormatVersion.first) + ", " + + std::to_string(OutputFormatVersion.second); + throw online_compile_error(std::string("The output format version (") + + Version + ") is not supported yet"); + } + + return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + DeviceStepping, CompileToSPIRVHandle, + FreeSPIRVOutputsHandle, UserArgs); +} + +template <> +template <> +std::vector online_compiler::compile( + const std::string &Source, const std::vector &UserArgs) { + + if (OutputFormatVersion != std::pair{0, 0}) { + std::string Version = std::to_string(OutputFormatVersion.first) + ", " + + std::to_string(OutputFormatVersion.second); + throw online_compile_error(std::string("The output format version (") + + Version + ") is not supported yet"); + } + + std::vector CMUserArgs = UserArgs; + CMUserArgs.push_back("-cmc"); + return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + DeviceStepping, CompileToSPIRVHandle, + FreeSPIRVOutputsHandle, CMUserArgs); +} + +} // namespace INTEL +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index e8677277fca0..452e39230f62 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3626,6 +3626,10 @@ _ZN2cl4sycl5eventC1Ev _ZN2cl4sycl5eventC2EP9_cl_eventRKNS0_7contextE _ZN2cl4sycl5eventC2ESt10shared_ptrINS0_6detail10event_implEE _ZN2cl4sycl5eventC2Ev +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJEEESt6vectorIhSaIhEERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEDpRKT_ +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJEEESt6vectorIhSaIhEERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEDpRKT_ _ZN2cl4sycl5queue10mem_adviseEPKvm14_pi_mem_advice _ZN2cl4sycl5queue10wait_proxyERKNS0_6detail13code_locationE _ZN2cl4sycl5queue11submit_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationE diff --git a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp new file mode 100644 index 000000000000..6dc0d09b3645 --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp @@ -0,0 +1,46 @@ +// REQUIRES: level_zero + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -lze_loader %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include +#include + +#include + +// clang-format off +#include +#include +// clang-format on + +using byte = unsigned char; + +sycl::kernel getSYCLKernelWithIL(sycl::context &Context, + const std::vector &IL) { + + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = IL.size(); + ZeModuleDesc.pInputModule = IL.data(); + ZeModuleDesc.pBuildFlags = ""; + ZeModuleDesc.pConstants = nullptr; + + assert(Context.get_devices().size() == 1 && "Expected to have only 1 device"); + sycl::device Device = Context.get_devices()[0]; + auto ZeDevice = Device.get_native(); + auto ZeContext = Context.get_native(); + + ze_module_build_log_handle_t ZeBuildLog; + ze_module_handle_t ZeModule; + ze_result_t ZeResult = zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, + &ZeModule, &ZeBuildLog); + if (ZeResult != ZE_RESULT_SUCCESS) + throw sycl::INTEL::online_compile_error(std::string("ZeResult = ") + + std::to_string(ZeResult)); + sycl::program SyclProgram = + sycl::level_zero::make(Context, ZeModule); + return SyclProgram.get_kernel("my_kernel"); +} + +#include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp new file mode 100644 index 000000000000..5db337d776d0 --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp @@ -0,0 +1,33 @@ +// REQUIRES: opencl + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -lOpenCL -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +#include +#include + +#include + +using byte = unsigned char; + +sycl::kernel getSYCLKernelWithIL(sycl::context &Context, + const std::vector &IL) { + cl_int Err; + cl_program ClProgram = + clCreateProgramWithIL(Context.get(), IL.data(), IL.size(), &Err); + if (Err != CL_SUCCESS) + throw sycl::compile_program_error(); + + Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr); + if (Err != CL_SUCCESS) + throw sycl::runtime_error(); + + cl_kernel ClKernel = clCreateKernel(ClProgram, "my_kernel", &Err); + if (Err != CL_SUCCESS) + throw sycl::runtime_error(); + + return sycl::kernel(ClKernel, Context); +} + +#include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp new file mode 100644 index 000000000000..4d2a310b10b0 --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -0,0 +1,171 @@ +#include +#include + +#include +#include + +static const char *CLSource = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0); + out[i] = in[i]*2 + 100; +} +)==="; + +static const char *CLSourceSyntaxError = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + syntax error here + size_t i = get_global_id(0); + out[i] = in[i]*2 + 100; +} +)==="; + +static const char *CMSource = R"===( +extern "C" +void cm_kernel() { +} +)==="; + +using namespace sycl::INTEL; + +void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { + std::cout << "Run the kernel now:\n"; + const int N = 4; + int InputArray[N] = {0, 1, 2, 3}; + int OutputArray[N] = {}; + + sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); + sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); + + Q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }).wait(); + + auto Out = OutputBuf.get_access(); + for (int I = 0; I < N; I++) { + std::cout << I << "*2 + 100 = " << Out[I] << "\n"; + } +} + +template +void doCompileAndRunTest(const std::string &Source) { + online_compiler Compiler; +} + +int main(int argc, char **argv) { + cl::sycl::queue Q; + cl::sycl::context Context = Q.get_context(); + cl::sycl::device Device = Q.get_device(); + + { // Compile and run a trivial OpenCL kernel. + std::cout << "Test case1\n"; + online_compiler Compiler; + std::vector IL; + try { + IL = Compiler.compile( + std::string(CLSource), + // Intentionally use one option twice. + std::vector{std::string("-cl-fast-relaxed-math"), + std::string("-cl-fast-relaxed-math")}); + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << std::string(e.what()) + << "\n"; + return 1; + } + testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); + } + + { // Compile and run a trivial OpenCL kernel using online_compiler() + // constructor accepting SYCL device. + std::cout << "Test case2\n"; + online_compiler Compiler(Device); + std::vector IL; + try { + IL = Compiler.compile(std::string(CLSource)); + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << std::string(e.what()) + << "\n"; + return 1; + } + testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); + } + + { // Compile a trivial CM kernel. + std::cout << "Test case3\n"; + online_compiler Compiler; + try { + std::vector IL = Compiler.compile(std::string(CMSource)); + + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << std::string(e.what()) + << "\n"; + return 1; + } + } + + { // Compile a source with syntax errors. + std::cout << "Test case4\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + IL = Compiler.compile(std::string(CLSourceSyntaxError)); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("syntax error here") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && "Failed to throw an exception for syntax error"); + } + + { // Compile a good CL source using unrecognized compilation options. + std::cout << "Test case5\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + IL = Compiler.compile( + std::string(CLSource), + // Intentionally use incorrect option. + std::vector{std::string("WRONG_OPTION")}); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("WRONG_OPTION") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && + "Failed to throw an exception for unrecognized option"); + } + + { // Try compiling CM source with OpenCL compiler. + std::cout << "Test case6\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + // Intentionally pass CMSource instead of CLSource. + IL = Compiler.compile(std::string(CMSource)); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("error: expected identifier or '('") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && "Failed to throw an exception for wrong program"); + } + + std::cout << "\nTest passed.\n"; + return 0; +} From 392cc468cb3b1d41a23c32b9546339c0b3c19057 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Mon, 21 Dec 2020 18:01:04 -0800 Subject: [PATCH 02/10] Address all reviewer's comments except 2: a) test on host; b) pull only ocloc_api.h from git Signed-off-by: Vyacheslav N Klochkov --- .../include/CL/sycl/INTEL/online_compiler.hpp | 6 +-- .../online_compiler/online_compiler.cpp | 45 +++++++++++++------ .../online_compiler/online_compiler_L0.cpp | 7 ++- .../online_compiler_OpenCL.cpp | 10 +++-- .../online_compiler_common.hpp | 20 ++++----- 5 files changed, 57 insertions(+), 31 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index bbef18c76dda..ae1a39905fe3 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -82,8 +82,8 @@ template class __SYCL_EXPORT online_compiler { /// device capabilities are supported by the target device(s). online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) : OutputFormat(fmt), OutputFormatVersion({0, 0}), - DeviceType(sycl::info::device_type::all), - DeviceArch(device_arch::any), Is64Bit(true), DeviceStepping("") {} + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler which targets given architecture and produces /// given compiled code format. Produces 64-bit device code. @@ -111,7 +111,7 @@ template class __SYCL_EXPORT online_compiler { /// can be different for different languages. /// Throws online_compile_error if compilation is not successful. template - std::vector compile(const std::string &src, const Tys &... args); + std::vector compile(const std::string &src, const Tys &...args); /// Sets the compiled code format of the compilation target and returns *this. online_compiler &setOutputFormat(compiled_code_format fmt) { diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp index e9187ab469c6..cbbadb89449f 100644 --- a/sycl/source/detail/online_compiler/online_compiler.cpp +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// #include -#include // for __SYCL_INLINE_NAMESPACE #include #include @@ -40,6 +39,8 @@ prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, Args.push_back("skl"); } } else { + // TODO: change that to generic device when ocloc adds support for it. + // For now "skl" is used as the lowest arch with GEN9 arch. Args.push_back("skl"); } @@ -58,6 +59,20 @@ prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, return Args; } +/// Compiles the given source \p Source to SPIR-V IL and returns IL as a vector +/// of bytes. +/// @param Source - Either OpenCL or CM source code. +/// @param DeviceType - SYCL device type, e.g. cpu, gpu, accelerator, etc. +/// @param DeviceArch - More detailed info on the target device architecture. +/// @param Is64Bit - If set to true, specifies the 64-bit architecture. +/// Otherwise, 32-bit is assumed. +/// @param DeviceStepping - implementation specific target device stepping. +/// @param CompileToSPIRVHandle - Output parameter. It is set to the address +/// of the library function doing the compilation. +/// @param FreeSPIRVOutputsHandle - Output parameter. It is set to the address +/// of the library function freeing memory +/// allocated during the compilation. +/// @param UserArgs - User's options to ocloc compiler. static std::vector compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, device_arch DeviceArch, bool Is64Bit, @@ -67,14 +82,14 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, if (!CompileToSPIRVHandle) { #ifdef __SYCL_RT_OS_WINDOWS - std::string OclocLibraryName = "ocloc64.dll"; + static const std::string OclocLibraryName = "ocloc64.dll"; #else - std::string OclocLibraryName = "libocloc.so"; + static const std::string OclocLibraryName = "libocloc.so"; #endif void *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); if (!OclocLibrary) - throw online_compile_error("Cannot load ocloc library"); - + throw online_compile_error("Cannot load ocloc library: " + + OclocLibraryName); void *OclocVersionHandle = sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion"); // The initial versions of ocloc library did not have the oclocVersion() @@ -132,7 +147,7 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, decltype(::oclocInvoke) *OclocInvokeFunc = reinterpret_cast(CompileToSPIRVHandle); - int Error = + int CompileError = OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths, &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs, &Outputs, &OutputLengths, &OutputNames); @@ -142,7 +157,8 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, size_t SpirVSize = 0; for (uint32_t I = 0; I < NumOutputs; I++) { size_t NameLen = strlen(OutputNames[I]); - if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr) { + if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr && + Outputs[I] != nullptr) { SpirVSize = OutputLengths[I]; SpirV = new byte[SpirVSize]; std::memcpy(SpirV, Outputs[I], SpirVSize); @@ -151,16 +167,19 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, } } - if (Error) - throw online_compile_error("ocloc reported compilation errors {" + - CompileLog + "}"); - + // Try to free memory before reporting possible error. decltype(::oclocFreeOutput) *OclocFreeOutputFunc = reinterpret_cast(FreeSPIRVOutputsHandle); - Error = + int MemFreeError = OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames); - if (Error) + if (CompileError) + throw online_compile_error("ocloc reported compilation errors: {\n" + + CompileLog + "\n}"); + if (!SpirV) + throw online_compile_error( + "Unexpected output: ocloc did not return SPIR-V"); + if (MemFreeError) throw online_compile_error("ocloc cannot safely free resources"); return std::vector(SpirV, SpirV + SpirVSize); diff --git a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp index 6dc0d09b3645..bf6cc0f8f240 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp @@ -4,6 +4,10 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// This test checks INTEL feature class online_compiler for Level-Zero. +// All Level-Zero specific code is kept here and the common part that can be +// re-used by other backends is kept in online_compiler_common.hpp file. + #include #include @@ -36,8 +40,7 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context, ze_result_t ZeResult = zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, &ZeModule, &ZeBuildLog); if (ZeResult != ZE_RESULT_SUCCESS) - throw sycl::INTEL::online_compile_error(std::string("ZeResult = ") + - std::to_string(ZeResult)); + throw sycl::runtime_error(); sycl::program SyclProgram = sycl::level_zero::make(Context, ZeModule); return SyclProgram.get_kernel("my_kernel"); diff --git a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp index 5db337d776d0..c8424a52f845 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp @@ -4,6 +4,10 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// This test checks INTEL feature class online_compiler for OpenCL. +// All OpenCL specific code is kept here and the common part that can be +// re-used by other backends is kept in online_compiler_common.hpp file. + #include #include @@ -17,15 +21,15 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context, cl_program ClProgram = clCreateProgramWithIL(Context.get(), IL.data(), IL.size(), &Err); if (Err != CL_SUCCESS) - throw sycl::compile_program_error(); + throw sycl::runtime_error("clCreateProgramWithIL() failed", Err); Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr); if (Err != CL_SUCCESS) - throw sycl::runtime_error(); + throw sycl::runtime_error("clBuildProgram() failed", Err); cl_kernel ClKernel = clCreateKernel(ClProgram, "my_kernel", &Err); if (Err != CL_SUCCESS) - throw sycl::runtime_error(); + throw sycl::runtime_error("clCreateKernel() failed", Err); return sycl::kernel(ClKernel, Context); } diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp index 4d2a310b10b0..e2518f56f7f6 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_common.hpp +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -43,14 +43,8 @@ void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { }).wait(); auto Out = OutputBuf.get_access(); - for (int I = 0; I < N; I++) { + for (int I = 0; I < N; I++) std::cout << I << "*2 + 100 = " << Out[I] << "\n"; - } -} - -template -void doCompileAndRunTest(const std::string &Source) { - online_compiler Compiler; } int main(int argc, char **argv) { @@ -65,9 +59,9 @@ int main(int argc, char **argv) { try { IL = Compiler.compile( std::string(CLSource), - // Intentionally use one option twice. + // Pass two options to check that more than one is accepted. std::vector{std::string("-cl-fast-relaxed-math"), - std::string("-cl-fast-relaxed-math")}); + std::string("-cl-finite-math-only")}); std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { @@ -125,6 +119,8 @@ int main(int argc, char **argv) { std::cerr << "Unexpected exception: " << Msg << "\n"; } assert(TestPassed && "Failed to throw an exception for syntax error"); + if (!TestPassed) + return 1; } { // Compile a good CL source using unrecognized compilation options. @@ -146,6 +142,8 @@ int main(int argc, char **argv) { } assert(TestPassed && "Failed to throw an exception for unrecognized option"); + if (!TestPassed) + return 1; } { // Try compiling CM source with OpenCL compiler. @@ -164,8 +162,10 @@ int main(int argc, char **argv) { std::cerr << "Unexpected exception: " << Msg << "\n"; } assert(TestPassed && "Failed to throw an exception for wrong program"); + if (!TestPassed) + return 1; } - std::cout << "\nTest passed.\n"; + std::cout << "\nAll test cases passed.\n"; return 0; } From df799bdad967ffc759d170f7a517a7bae88017af Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Tue, 22 Dec 2020 09:38:36 -0800 Subject: [PATCH 03/10] Fix the libocloc library version check Signed-off-by: Vyacheslav N Klochkov --- sycl/include/CL/sycl/INTEL/online_compiler.hpp | 2 +- .../detail/online_compiler/online_compiler.cpp | 14 ++++++++++---- 2 files changed, 11 insertions(+), 5 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index ae1a39905fe3..275451b6afa8 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -111,7 +111,7 @@ template class __SYCL_EXPORT online_compiler { /// can be different for different languages. /// Throws online_compile_error if compilation is not successful. template - std::vector compile(const std::string &src, const Tys &...args); + std::vector compile(const std::string &src, const Tys &... args); /// Sets the compiled code format of the compilation target and returns *this. online_compiler &setOutputFormat(compiled_code_format fmt) { diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp index cbbadb89449f..e2fc8feda765 100644 --- a/sycl/source/detail/online_compiler/online_compiler.cpp +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -101,15 +101,21 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, reinterpret_cast(OclocVersionHandle); LoadedVersion = OclocVersionFunc(); } + // The loaded library with version (A.B) is compatible with expected API/ABI + // version (X.Y) used here if A == B and B >= Y. int LoadedVersionMajor = LoadedVersion >> 16; + int LoadedVersionMinor = LoadedVersion & 0xffff; int CurrentVersionMajor = (ocloc_version_t::OCLOC_VERSION_CURRENT) >> 16; - if (LoadedVersionMajor != CurrentVersionMajor) + int CurrentVersionMinor = (ocloc_version_t::OCLOC_VERSION_CURRENT) & 0xffff; + if (LoadedVersionMajor != CurrentVersionMajor || + LoadedVersionMinor < CurrentVersionMinor) throw online_compile_error( std::string("Found incompatible version of ocloc library: (") + - std::to_string(LoadedVersionMajor) + ", " + - std::to_string(LoadedVersion & 0xffff) + + std::to_string(LoadedVersionMajor) + "." + + std::to_string(LoadedVersionMinor) + "). The supported versions are (" + - std::to_string(CurrentVersionMajor) + ", *)."); + std::to_string(CurrentVersionMajor) + ". N), where (N >= " + + std::to_string(CurrentVersionMinor) + ")."); CompileToSPIRVHandle = sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); From 1d5671dd73c485d6ed1497d63eff567dbf8d45f2 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Tue, 22 Dec 2020 11:21:30 -0800 Subject: [PATCH 04/10] Fix clang-format and fix linkage error on Widows Signed-off-by: Vyacheslav N Klochkov --- sycl/include/CL/sycl/INTEL/online_compiler.hpp | 7 ++++--- .../detail/online_compiler/online_compiler.cpp | 13 +++++++------ 2 files changed, 11 insertions(+), 9 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index 275451b6afa8..fd87eb2cde8b 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -74,7 +74,7 @@ enum class source_language { opencl_c, cm }; /// Represents an online compiler for the language given as template /// parameter. -template class __SYCL_EXPORT online_compiler { +template class online_compiler { public: /// Constructs online compiler which can target any device and produces /// given compiled code format. Produces 64-bit device code. @@ -190,7 +190,8 @@ template class __SYCL_EXPORT online_compiler { /// OpenCL JIT compiler options must be supported. template <> template <> -std::vector online_compiler::compile( +__SYCL_EXPORT std::vector +online_compiler::compile( const std::string &src, const std::vector &options); /// Compiles the given OpenCL source. May throw \c online_compile_error. @@ -207,7 +208,7 @@ online_compiler::compile(const std::string &src) { /// @param options - compilation options (implementation defined). template <> template <> -std::vector online_compiler::compile( +__SYCL_EXPORT std::vector online_compiler::compile( const std::string &src, const std::vector &options); /// Compiles the given CM source \p src. diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp index e2fc8feda765..ca352c387b8a 100644 --- a/sycl/source/detail/online_compiler/online_compiler.cpp +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -105,8 +105,8 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, // version (X.Y) used here if A == B and B >= Y. int LoadedVersionMajor = LoadedVersion >> 16; int LoadedVersionMinor = LoadedVersion & 0xffff; - int CurrentVersionMajor = (ocloc_version_t::OCLOC_VERSION_CURRENT) >> 16; - int CurrentVersionMinor = (ocloc_version_t::OCLOC_VERSION_CURRENT) & 0xffff; + int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16; + int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff; if (LoadedVersionMajor != CurrentVersionMajor || LoadedVersionMinor < CurrentVersionMinor) throw online_compile_error( @@ -114,8 +114,8 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, std::to_string(LoadedVersionMajor) + "." + std::to_string(LoadedVersionMinor) + "). The supported versions are (" + - std::to_string(CurrentVersionMajor) + ". N), where (N >= " + - std::to_string(CurrentVersionMinor) + ")."); + std::to_string(CurrentVersionMajor) + + ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ")."); CompileToSPIRVHandle = sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); @@ -194,7 +194,8 @@ compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, template <> template <> -std::vector online_compiler::compile( +__SYCL_EXPORT std::vector +online_compiler::compile( const std::string &Source, const std::vector &UserArgs) { if (OutputFormatVersion != std::pair{0, 0}) { @@ -211,7 +212,7 @@ std::vector online_compiler::compile( template <> template <> -std::vector online_compiler::compile( +__SYCL_EXPORT std::vector online_compiler::compile( const std::string &Source, const std::vector &UserArgs) { if (OutputFormatVersion != std::pair{0, 0}) { From e73322c2e99efe880b52bd36ffd8c9a62cc3d429 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Tue, 22 Dec 2020 13:18:38 -0800 Subject: [PATCH 05/10] Run the LIT tests on HOST as well Signed-off-by: Vyacheslav N Klochkov --- .../on-device/online_compiler/online_compiler_L0.cpp | 6 +++++- .../online_compiler/online_compiler_OpenCL.cpp | 6 +++++- .../online_compiler/online_compiler_common.hpp | 10 ++++++++++ 3 files changed, 20 insertions(+), 2 deletions(-) diff --git a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp index bf6cc0f8f240..e74437e4b1ce 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp @@ -1,8 +1,10 @@ // REQUIRES: level_zero -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -lze_loader %s -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -DRUN_KERNELS -lze_loader %s -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -lze_loader %s -o %th.out +// RUN: %RUN_ON_HOST %th.out // This test checks INTEL feature class online_compiler for Level-Zero. // All Level-Zero specific code is kept here and the common part that can be @@ -20,6 +22,7 @@ using byte = unsigned char; +#ifdef RUN_KERNELS sycl::kernel getSYCLKernelWithIL(sycl::context &Context, const std::vector &IL) { @@ -45,5 +48,6 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context, sycl::level_zero::make(Context, ZeModule); return SyclProgram.get_kernel("my_kernel"); } +#endif // RUN_KERNELS #include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp index c8424a52f845..7aff6e6037ba 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp @@ -1,8 +1,10 @@ // REQUIRES: opencl -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -lOpenCL -o %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -DRUN_KERNELS -lOpenCL -o %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -lOpenCL -o %th.out +// RUN: %RUN_ON_HOST %th.out // This test checks INTEL feature class online_compiler for OpenCL. // All OpenCL specific code is kept here and the common part that can be @@ -15,6 +17,7 @@ using byte = unsigned char; +#ifdef RUN_KERNELS sycl::kernel getSYCLKernelWithIL(sycl::context &Context, const std::vector &IL) { cl_int Err; @@ -33,5 +36,6 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context, return sycl::kernel(ClKernel, Context); } +#endif // RUN_KERNELS #include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp index e2518f56f7f6..a8c4fbe66756 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_common.hpp +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -27,6 +27,7 @@ void cm_kernel() { using namespace sycl::INTEL; +#ifdef RUN_KERNELS void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { std::cout << "Run the kernel now:\n"; const int N = 4; @@ -46,6 +47,7 @@ void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { for (int I = 0; I < N; I++) std::cout << I << "*2 + 100 = " << Out[I] << "\n"; } +#endif // RUN_KERNELS int main(int argc, char **argv) { cl::sycl::queue Q; @@ -69,7 +71,9 @@ int main(int argc, char **argv) { << "\n"; return 1; } +#ifdef RUN_KERNELS testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); +#endif // RUN_KERNELS } { // Compile and run a trivial OpenCL kernel using online_compiler() @@ -86,9 +90,14 @@ int main(int argc, char **argv) { << "\n"; return 1; } +#ifdef RUN_KERNELS testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); +#endif // RUN_KERNELS } +#ifdef COMPILE_CM_KERNEL + // TODO: this test is temporarily turned off because CI buildbots do not set + // PATHs to clangFEWrapper library properly. { // Compile a trivial CM kernel. std::cout << "Test case3\n"; online_compiler Compiler; @@ -103,6 +112,7 @@ int main(int argc, char **argv) { return 1; } } +#endif // COMPILE_CM_KERNEL { // Compile a source with syntax errors. std::cout << "Test case4\n"; From 893852b81738ba82be3748c1cfcbabc8f61bd94d Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Tue, 22 Dec 2020 23:12:27 -0800 Subject: [PATCH 06/10] Address the review comments from keryell Signed-off-by: Vyacheslav N Klochkov --- .../online_compiler_common.hpp | 37 +++++++++---------- 1 file changed, 18 insertions(+), 19 deletions(-) diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp index a8c4fbe66756..cee2e5267e51 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_common.hpp +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -4,14 +4,14 @@ #include #include -static const char *CLSource = R"===( +static constexpr char *CLSource = R"===( __kernel void my_kernel(__global int *in, __global int *out) { size_t i = get_global_id(0); out[i] = in[i]*2 + 100; } )==="; -static const char *CLSourceSyntaxError = R"===( +static constexpr char *CLSourceSyntaxError = R"===( __kernel void my_kernel(__global int *in, __global int *out) { syntax error here size_t i = get_global_id(0); @@ -19,7 +19,7 @@ __kernel void my_kernel(__global int *in, __global int *out) { } )==="; -static const char *CMSource = R"===( +static constexpr char *CMSource = R"===( extern "C" void cm_kernel() { } @@ -30,7 +30,7 @@ using namespace sycl::INTEL; #ifdef RUN_KERNELS void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { std::cout << "Run the kernel now:\n"; - const int N = 4; + constexpr int N = 4; int InputArray[N] = {0, 1, 2, 3}; int OutputArray[N] = {}; @@ -38,10 +38,10 @@ void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); Q.submit([&](sycl::handler &CGH) { - CGH.set_arg(0, InputBuf.get_access(CGH)); - CGH.set_arg(1, OutputBuf.get_access(CGH)); - CGH.parallel_for(sycl::range<1>{N}, Kernel); - }).wait(); + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }); auto Out = OutputBuf.get_access(); for (int I = 0; I < N; I++) @@ -60,10 +60,10 @@ int main(int argc, char **argv) { std::vector IL; try { IL = Compiler.compile( - std::string(CLSource), + CLSource, // Pass two options to check that more than one is accepted. - std::vector{std::string("-cl-fast-relaxed-math"), - std::string("-cl-finite-math-only")}); + std::vector{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}); std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { @@ -82,7 +82,7 @@ int main(int argc, char **argv) { online_compiler Compiler(Device); std::vector IL; try { - IL = Compiler.compile(std::string(CLSource)); + IL = Compiler.compile(CLSource); std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { @@ -102,7 +102,7 @@ int main(int argc, char **argv) { std::cout << "Test case3\n"; online_compiler Compiler; try { - std::vector IL = Compiler.compile(std::string(CMSource)); + std::vector IL = Compiler.compile(CMSource); std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); @@ -120,7 +120,7 @@ int main(int argc, char **argv) { std::vector IL; bool TestPassed = false; try { - IL = Compiler.compile(std::string(CLSourceSyntaxError)); + IL = Compiler.compile(CLSourceSyntaxError); } catch (sycl::exception &e) { std::string Msg = e.what(); if (Msg.find("syntax error here") != std::string::npos) @@ -139,10 +139,9 @@ int main(int argc, char **argv) { std::vector IL; bool TestPassed = false; try { - IL = Compiler.compile( - std::string(CLSource), - // Intentionally use incorrect option. - std::vector{std::string("WRONG_OPTION")}); + IL = Compiler.compile(CLSource, + // Intentionally use incorrect option. + std::vector{"WRONG_OPTION"}); } catch (sycl::exception &e) { std::string Msg = e.what(); if (Msg.find("WRONG_OPTION") != std::string::npos) @@ -163,7 +162,7 @@ int main(int argc, char **argv) { bool TestPassed = false; try { // Intentionally pass CMSource instead of CLSource. - IL = Compiler.compile(std::string(CMSource)); + IL = Compiler.compile(CMSource); } catch (sycl::exception &e) { std::string Msg = e.what(); if (Msg.find("error: expected identifier or '('") != std::string::npos) From c1382afa78cdce5ee97c7eaa23d8f9fea3478a8b Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Wed, 23 Dec 2020 08:31:08 -0800 Subject: [PATCH 07/10] More reviewer's comments addressed Signed-off-by: Vyacheslav N Klochkov --- .../online_compiler/online_compiler_common.hpp | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp index cee2e5267e51..9ef97c45020b 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_common.hpp +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -4,14 +4,14 @@ #include #include -static constexpr char *CLSource = R"===( +auto constexpr CLSource = R"===( __kernel void my_kernel(__global int *in, __global int *out) { size_t i = get_global_id(0); out[i] = in[i]*2 + 100; } )==="; -static constexpr char *CLSourceSyntaxError = R"===( +auto constexpr CLSourceSyntaxError = R"===( __kernel void my_kernel(__global int *in, __global int *out) { syntax error here size_t i = get_global_id(0); @@ -19,7 +19,7 @@ __kernel void my_kernel(__global int *in, __global int *out) { } )==="; -static constexpr char *CMSource = R"===( +auto constexpr CMSource = R"===( extern "C" void cm_kernel() { } @@ -67,8 +67,7 @@ int main(int argc, char **argv) { std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { - std::cout << "Compilation to IL failed: " << std::string(e.what()) - << "\n"; + std::cout << "Compilation to IL failed: " << e.what() << "\n"; return 1; } #ifdef RUN_KERNELS @@ -86,8 +85,7 @@ int main(int argc, char **argv) { std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { - std::cout << "Compilation to IL failed: " << std::string(e.what()) - << "\n"; + std::cout << "Compilation to IL failed: " << e.what() << "\n"; return 1; } #ifdef RUN_KERNELS @@ -107,8 +105,7 @@ int main(int argc, char **argv) { std::cout << "IL size = " << IL.size() << "\n"; assert(IL.size() > 0 && "Unexpected IL size"); } catch (sycl::exception &e) { - std::cout << "Compilation to IL failed: " << std::string(e.what()) - << "\n"; + std::cout << "Compilation to IL failed: " << e.what() << "\n"; return 1; } } From 2fd5d6f62a46f13131a344b1c601f232cad9f64c Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Thu, 24 Dec 2020 12:25:54 -0800 Subject: [PATCH 08/10] Update the SYCL ABI test after all those additional fixes in this PR Signed-off-by: Vyacheslav N Klochkov --- sycl/test/abi/sycl_symbols_linux.dump | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 452e39230f62..27c2a4082eb2 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3627,9 +3627,7 @@ _ZN2cl4sycl5eventC2EP9_cl_eventRKNS0_7contextE _ZN2cl4sycl5eventC2ESt10shared_ptrINS0_6detail10event_implEE _ZN2cl4sycl5eventC2Ev _ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ -_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJEEESt6vectorIhSaIhEERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEDpRKT_ _ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ -_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJEEESt6vectorIhSaIhEERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEDpRKT_ _ZN2cl4sycl5queue10mem_adviseEPKvm14_pi_mem_advice _ZN2cl4sycl5queue10wait_proxyERKNS0_6detail13code_locationE _ZN2cl4sycl5queue11submit_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationE From f555c9767d24bfe3989d134ed08cfd4fda70494a Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Thu, 24 Dec 2020 13:08:21 -0800 Subject: [PATCH 09/10] Additional fixes in online_compile_error class after merging conflicts with exception.hpp Signed-off-by: Vyacheslav N Klochkov --- sycl/include/CL/sycl/INTEL/online_compiler.hpp | 7 +------ sycl/include/CL/sycl/exception.hpp | 2 ++ 2 files changed, 3 insertions(+), 6 deletions(-) diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index fd87eb2cde8b..f8bb63bc79dc 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -61,12 +61,7 @@ class device_arch { class online_compile_error : public sycl::exception { public: online_compile_error() = default; - online_compile_error(const string_class &Msg) : Msg(Msg) {} - - const char *what() const noexcept override { return Msg.c_str(); }; - -private: - string_class Msg; + online_compile_error(const string_class &Msg) : sycl::exception(Msg) {} }; /// Designates a source language for the online compiler. diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp index 038ddeb94fc7..10c0d55a47d0 100644 --- a/sycl/include/CL/sycl/exception.hpp +++ b/sycl/include/CL/sycl/exception.hpp @@ -52,6 +52,8 @@ class __SYCL_EXPORT exception : public std::exception { shared_ptr_class Context = nullptr) : MMsg(Msg + " " + detail::codeToString(CLErr)), MCLErr(CLErr), MContext(Context) {} + + exception(const string_class &Msg) : MMsg(Msg), MContext(nullptr) {} }; class runtime_error : public exception { From ee66600cea8b23e73ac2874b8c364fada556c805 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Thu, 24 Dec 2020 13:15:54 -0800 Subject: [PATCH 10/10] Enable the CM compilation test case for Windows only (temporarily turned off on Linux) Signed-off-by: Vyacheslav N Klochkov --- sycl/test/on-device/online_compiler/online_compiler_L0.cpp | 7 +++++++ .../on-device/online_compiler/online_compiler_OpenCL.cpp | 7 +++++++ 2 files changed, 14 insertions(+) diff --git a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp index e74437e4b1ce..5fd556f81a4d 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp @@ -15,6 +15,13 @@ #include +// TODO: The testing of CM compilation is temporarily turned OFF on Linux +// due to problems with dependencies on libclangFEWrapper.so which is not +// currently included into NEO package on Linux. +#ifdef _WIN32 +#define COMPILE_CM_KERNEL 1 +#endif + // clang-format off #include #include diff --git a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp index 7aff6e6037ba..086b2da992f1 100644 --- a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp +++ b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp @@ -10,6 +10,13 @@ // All OpenCL specific code is kept here and the common part that can be // re-used by other backends is kept in online_compiler_common.hpp file. +// TODO: The testing of CM compilation is temporarily turned OFF on Linux +// due to problems with dependencies on libclangFEWrapper.so which is not +// currently included into NEO package on Linux. +#ifdef _WIN32 +#define COMPILE_CM_KERNEL 1 +#endif + #include #include