From cd1acf3b98b8298f733293e627a821a8b35e158b Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Sun, 30 May 2021 21:26:18 +0300 Subject: [PATCH 1/2] [SYCL] Fix build options merge in program manager The program manager gets build options from several sources including device code image, SYCL API and environment variables. Make sure that options are space separated. --- .../program_manager/program_manager.cpp | 120 ++++---- .../get_native_interop/test_get_native.cpp | 2 + sycl/unittests/helpers/PiImage.hpp | 20 +- sycl/unittests/misc/CMakeLists.txt | 3 +- sycl/unittests/misc/KernelBuildOptions.cpp | 277 ++++++++++++++++++ 5 files changed, 355 insertions(+), 67 deletions(-) create mode 100644 sycl/unittests/misc/KernelBuildOptions.cpp diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index e0973cc3b1f16..3e243b83b9963 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -347,6 +347,54 @@ RT::PiProgram ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img, return Res; } +static void applyOptionsFromImage(std::string &CompileOpts, + std::string &LinkOpts, + const RTDeviceBinaryImage &Img) { + // Build options are overridden if environment variables are present. + // Environment variables are not changed during program lifecycle so it + // is reasonable to use static here to read them only once. + static const char *CompileOptsEnv = + SYCLConfig::get(); + static const char *LinkOptsEnv = SYCLConfig::get(); + // Update only if compile options are not overwritten by environment + // variable + if (!CompileOptsEnv) { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += Img.getCompileOptions(); + } + + // The -vc-codegen option is always preserved for ESIMD kernels, regardless + // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. + pi_device_binary_property isEsimdImage = Img.getProperty("isEsimdImage"); + if (isEsimdImage && pi::DeviceBinaryProperty(isEsimdImage).asUint32()) { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += "-vc-codegen"; + } + + // Update only if link options are not overwritten by environment variable + if (!LinkOptsEnv) + if (!LinkOpts.empty()) + LinkOpts += " "; + LinkOpts += Img.getLinkOptions(); +} + +static void applyOptionsFromEnvironment(std::string &CompileOpts, + std::string &LinkOpts) { + // Build options are overridden if environment variables are present. + // Environment variables are not changed during program lifecycle so it + // is reasonable to use static here to read them only once. + static const char *CompileOptsEnv = + SYCLConfig::get(); + if (CompileOptsEnv) { + CompileOpts = CompileOptsEnv; + } + static const char *LinkOptsEnv = SYCLConfig::get(); + if (LinkOptsEnv) { + LinkOpts = LinkOptsEnv; + } +} RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, @@ -374,26 +422,12 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M, std::string CompileOpts; std::string LinkOpts; - // Build options are overridden if environment variables are present. - // Environment variables are not changed during program lifecycle so it - // is reasonable to use static here to read them only once. - static const char *CompileOptsEnv = - SYCLConfig::get(); - if (CompileOptsEnv) { - CompileOpts = CompileOptsEnv; - } else { // Use build options only when the environment variable is missed - if (Prg) { - std::string BuildOptions = Prg->get_build_options(); - if (!BuildOptions.empty()) { - CompileOpts += " "; - CompileOpts += BuildOptions; - } - } - } - static const char *LinkOptsEnv = SYCLConfig::get(); - if (LinkOptsEnv) { - LinkOpts = LinkOptsEnv; + if (Prg) { + CompileOpts = Prg->get_build_options(); } + + applyOptionsFromEnvironment(CompileOpts, LinkOpts); + SerializedObj SpecConsts; if (Prg) Prg->stableSerializeSpecConstRegistry(SpecConsts); @@ -402,24 +436,8 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M, &LinkOpts, &JITCompilationIsRequired, SpecConsts] { const RTDeviceBinaryImage &Img = getDeviceImage(M, KSId, Context, Device, JITCompilationIsRequired); - // Update only if compile options are not overwritten by environment - // variable - if (!CompileOptsEnv) { - CompileOpts += Img.getCompileOptions(); - } - - // The -vc-codegen option is always preserved for ESIMD kernels, regardless - // of the contents SYCL_PROGRAM_COMPILE_OPTIONS environment variable. - pi_device_binary_property isEsimdImage = Img.getProperty("isEsimdImage"); - if (isEsimdImage && pi::DeviceBinaryProperty(isEsimdImage).asUint32()) { - if (!CompileOpts.empty()) - CompileOpts += " "; - CompileOpts += "-vc-codegen"; - } - // Update only if link options are not overwritten by environment variable - if (!LinkOptsEnv) - LinkOpts += Img.getLinkOptions(); + applyOptionsFromImage(CompileOpts, LinkOpts, Img); ContextImplPtr ContextImpl = getSyclObjImpl(Context); const detail::plugin &Plugin = ContextImpl->getPlugin(); RT::PiProgram NativePrg; @@ -1516,18 +1534,7 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, std::string CompileOpts; std::string LinkOpts; - // Build options are overridden if environment variables are present. - // Environment variables are not changed during program lifecycle so it - // is reasonable to use static here to read them only once. - static const char *CompileOptsEnv = - SYCLConfig::get(); - if (CompileOptsEnv) - CompileOpts = CompileOptsEnv; - - static const char *LinkOptsEnv = SYCLConfig::get(); - if (LinkOptsEnv) { - LinkOpts = LinkOptsEnv; - } + applyOptionsFromEnvironment(CompileOpts, LinkOpts); const RTDeviceBinaryImage *ImgPtr = InputImpl->get_bin_image_ref(); const RTDeviceBinaryImage &Img = *ImgPtr; @@ -1535,22 +1542,7 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, // TODO: Unify this code with getBuiltPIProgram auto BuildF = [this, &Context, Img, &Devs, &CompileOpts, &LinkOpts, &InputImpl] { - // Update only if compile options are not overwritten by environment - // variable - if (!CompileOptsEnv) { - CompileOpts += Img.getCompileOptions(); - pi_device_binary_property isEsimdImage = Img.getProperty("isEsimdImage"); - - if (isEsimdImage && pi::DeviceBinaryProperty(isEsimdImage).asUint32()) { - if (!CompileOpts.empty()) - CompileOpts += " "; - CompileOpts += "-vc-codegen"; - } - } - - // Update only if link options are not overwritten by environment variable - if (!LinkOptsEnv) - LinkOpts += Img.getLinkOptions(); + applyOptionsFromImage(CompileOpts, LinkOpts, Img); ContextImplPtr ContextImpl = getSyclObjImpl(Context); const detail::plugin &Plugin = ContextImpl->getPlugin(); diff --git a/sycl/unittests/get_native_interop/test_get_native.cpp b/sycl/unittests/get_native_interop/test_get_native.cpp index 0930011ff74e6..41cb3823d00a9 100644 --- a/sycl/unittests/get_native_interop/test_get_native.cpp +++ b/sycl/unittests/get_native_interop/test_get_native.cpp @@ -8,6 +8,8 @@ // //===----------------------------------------------------------------------===// +#define SYCL2020_DISABLE_DEPRECATION_WARNINGS + #include #include #include diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp index 6306f9961b28f..eee1fa4b20b6f 100644 --- a/sycl/unittests/helpers/PiImage.hpp +++ b/sycl/unittests/helpers/PiImage.hpp @@ -51,8 +51,13 @@ class PiProperty { private: void updateNativeType() { - MNative = NativeType{const_cast(MName.c_str()), - const_cast(MData.data()), MType, MData.size()}; + if ( MType == PI_PROPERTY_TYPE_UINT32 ) { + MNative = NativeType{const_cast(MName.c_str()), + nullptr, MType, *((uint32_t*)MData.data())}; + } else { + MNative = NativeType{const_cast(MName.c_str()), + const_cast(MData.data()), MType, MData.size()}; + } } std::string MName; std::vector MData; @@ -366,6 +371,17 @@ void addSpecConstants(PiArray SpecConstants, std::move(DefaultValues)); } +/// Utility function to add ESIMD kernel flag to property set. +void addESIMDFlag(PiPropertySet &Props) { + std::vector ValData(sizeof(uint32_t)); + ValData[0]= 1; + PiProperty Prop{"isEsimdImage", ValData, PI_PROPERTY_TYPE_UINT32}; + + PiArray Value{std::move(Prop)}; + + Props.insert(__SYCL_PI_PROPERTY_SET_SYCL_MISC_PROP, std::move(Value)); +} + /// Utility function to generate offload entries for kernels without arguments. PiArray makeEmptyKernels(std::initializer_list KernelNames) { diff --git a/sycl/unittests/misc/CMakeLists.txt b/sycl/unittests/misc/CMakeLists.txt index 169a1572b2d09..127089d85ebaa 100644 --- a/sycl/unittests/misc/CMakeLists.txt +++ b/sycl/unittests/misc/CMakeLists.txt @@ -1,6 +1,7 @@ set(sycl_lib_dir $) add_definitions(-DSYCL_LIB_DIR="${sycl_lib_dir}") add_sycl_unittest(MiscTests SHARED - OsUtils.cpp CircularBuffer.cpp + KernelBuildOptions.cpp + OsUtils.cpp ) diff --git a/sycl/unittests/misc/KernelBuildOptions.cpp b/sycl/unittests/misc/KernelBuildOptions.cpp new file mode 100644 index 0000000000000..b763ca1e6efb1 --- /dev/null +++ b/sycl/unittests/misc/KernelBuildOptions.cpp @@ -0,0 +1,277 @@ +//==- KernelBuildOptions.cpp - Kernel build options processing unit test --==// +// +// 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 +// +//===---------------------------------------------------------------------===// + +#define SYCL2020_DISABLE_DEPRECATION_WARNINGS + +#include +#include +#include + +#include + +class TestKernel; + +static std::string BuildOpts; +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +template <> struct KernelInfo { + static constexpr unsigned getNumParams() { return 0; } + static const kernel_param_desc_t &getParamDesc(int) { + static kernel_param_desc_t Dummy; + return Dummy; + } + static constexpr const char *getName() { return "TestKernel"; } + static constexpr bool isESIMD() { return true; } + static constexpr bool callsThisItem() { return false; } + static constexpr bool callsAnyThisFreeFunction() { return false; } +}; + +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + +static pi_result redefinedProgramCreate(pi_context, const void *, size_t, + pi_program *) { + return PI_SUCCESS; +} + +static pi_result redefinedProgramBuild( + pi_program prog, pi_uint32, const pi_device *, const char * options, + void (*pfn_notify)(pi_program program, void *user_data), void *user_data) { + if (options) + BuildOpts = options; + else + BuildOpts = ""; + if (pfn_notify) { + pfn_notify(prog, user_data); + } + return PI_SUCCESS; +} + +static pi_result redefinedProgramCompile(pi_program, pi_uint32, + const pi_device *, const char * options, + pi_uint32, const pi_program *, + const char **, + void (*)(pi_program, void *), void *) { + if (options) + BuildOpts = options; + else + BuildOpts = ""; + return PI_SUCCESS; +} + +static pi_result redefinedProgramLink(pi_context, pi_uint32, const pi_device *, + const char * options, pi_uint32, + const pi_program *, + void (*)(pi_program, void *), void *, + pi_program *) { + if (options) + BuildOpts = options; + else + BuildOpts = ""; + return PI_SUCCESS; +} + +static pi_result redefinedProgramGetInfo(pi_program program, + pi_program_info param_name, + size_t param_value_size, + void *param_value, + size_t *param_value_size_ret) { + if (param_name == PI_PROGRAM_INFO_NUM_DEVICES) { + auto value = reinterpret_cast(param_value); + *value = 1; + } + + if (param_name == PI_PROGRAM_INFO_BINARY_SIZES) { + auto value = reinterpret_cast(param_value); + value[0] = 1; + } + + if (param_name == PI_PROGRAM_INFO_BINARIES) { + auto value = reinterpret_cast(param_value); + value[0] = 1; + } + + return PI_SUCCESS; +} + +static pi_result redefinedProgramRetain(pi_program program) { + return PI_SUCCESS; +} + +static pi_result redefinedProgramRelease(pi_program program) { + return PI_SUCCESS; +} + +static pi_result redefinedKernelCreate(pi_program program, + const char *kernel_name, + pi_kernel *ret_kernel) { + *ret_kernel = reinterpret_cast(new int[1]); + return PI_SUCCESS; +} + +static pi_result redefinedKernelRetain(pi_kernel kernel) { return PI_SUCCESS; } + +static pi_result redefinedKernelRelease(pi_kernel kernel) { + delete[] reinterpret_cast(kernel); + return PI_SUCCESS; +} + +static pi_result redefinedKernelGetInfo(pi_kernel kernel, + pi_kernel_info param_name, + size_t param_value_size, + void *param_value, + size_t *param_value_size_ret) { + return PI_SUCCESS; +} + +static pi_result redefinedProgramCreateWithSource(pi_context context, + pi_uint32 count, + const char **strings, + const size_t *lengths, + pi_program *ret_program) { + return PI_SUCCESS; +} + +static pi_result redefinedProgramCreateWithBinary( + pi_context context, pi_uint32 num_devices, const pi_device *device_list, + const size_t *lengths, const unsigned char **binaries, + pi_int32 *binary_status, pi_program *ret_program) { + return PI_SUCCESS; +} + +static pi_result redefinedKernelSetExecInfo(pi_kernel kernel, + pi_kernel_exec_info value_name, + size_t param_value_size, + const void *param_value) { + return PI_SUCCESS; +} + +static pi_result redefinedEventsWait(pi_uint32 num_events, + const pi_event *event_list) { + return PI_SUCCESS; +} + +static pi_result redefinedEnqueueKernelLaunch(pi_queue, pi_kernel, pi_uint32, + const size_t *, const size_t *, + const size_t *, pi_uint32, + const pi_event *, pi_event *) { + return PI_SUCCESS; +} + +static void setupDefaultMockAPIs(sycl::unittest::PiMock &Mock) { + using namespace sycl::detail; + Mock.redefine( + redefinedProgramCreateWithSource); + Mock.redefine( + redefinedProgramCreateWithBinary); + Mock.redefine(redefinedProgramCreate); + Mock.redefine(redefinedProgramCompile); + Mock.redefine(redefinedProgramLink); + Mock.redefine(redefinedProgramBuild); + Mock.redefine(redefinedProgramGetInfo); + Mock.redefine(redefinedProgramRetain); + Mock.redefine(redefinedProgramRelease); + Mock.redefine(redefinedKernelCreate); + Mock.redefine(redefinedKernelRetain); + Mock.redefine(redefinedKernelRelease); + Mock.redefine(redefinedKernelGetInfo); + Mock.redefine(redefinedKernelSetExecInfo); + Mock.redefine(redefinedEventsWait); + Mock.redefine(redefinedEnqueueKernelLaunch); +} + +static sycl::unittest::PiImage generateDefaultImage() { + using namespace sycl::unittest; + + PiPropertySet PropSet; + addESIMDFlag(PropSet); + std::vector Bin{0, 1, 2, 3, 4, 5}; // Random data + + PiArray Entries = makeEmptyKernels({"TestKernel"}); + + PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV, // Format + __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec + "-compile-img", // Compile options + "-link-img", // Link options + std::move(Bin), + std::move(Entries), + std::move(PropSet)}; + + return Img; +} + +sycl::unittest::PiImage Img = generateDefaultImage(); +sycl::unittest::PiImageArray ImgArray{Img}; + +TEST(KernelBuildOptions, KernelBundleBasic) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + + sycl::queue Queue{Dev}; + + const sycl::context Ctx = Queue.get_context(); + + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle(Ctx, {Dev}); + auto ExecBundle = sycl::build(KernelBundle); + EXPECT_EQ(BuildOpts, "-compile-img -vc-codegen"); + + auto ObjBundle = sycl::compile(KernelBundle, KernelBundle.get_devices()); + // TODO: uncomment when image options are passed to BE + // EXPECT_EQ(BuildOpts, "-compile-img -vc-codegen"); + + auto LinkBundle = sycl::link(ObjBundle, ObjBundle.get_devices()); + // TODO: uncomment when image options are passed to BE + // EXPECT_EQ(BuildOpts, "-link-img -vc-codegen"); + +} + +TEST(KernelBuildOptions, Program) { + sycl::platform Plt{sycl::default_selector()}; + if (Plt.is_host()) { + std::cerr << "Test is not supported on host, skipping\n"; + return; // test is not supported on host. + } + + if (Plt.get_backend() == sycl::backend::cuda) { + std::cerr << "Test is not supported on CUDA platform, skipping\n"; + return; + } + + sycl::unittest::PiMock Mock{Plt}; + setupDefaultMockAPIs(Mock); + + const sycl::device Dev = Plt.get_devices()[0]; + + sycl::queue Queue{Dev}; + + const sycl::context Ctx = Queue.get_context(); + sycl::program Prg1(Ctx); + sycl::program Prg2(Ctx); + + Prg1.build_with_source(""); + EXPECT_TRUE(BuildOpts.size()==0) << "Expect empty build options"; + Prg2.build_with_source("", "-api-opts"); + EXPECT_EQ(BuildOpts, "-api-opts"); + +} From 38d221663204b6f5d18e9c2ba93aba95eb09ae50 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Mon, 31 May 2021 07:33:18 +0300 Subject: [PATCH 2/2] Fix clang-format issue --- sycl/unittests/helpers/PiImage.hpp | 13 +++++++------ sycl/unittests/misc/KernelBuildOptions.cpp | 14 ++++++-------- 2 files changed, 13 insertions(+), 14 deletions(-) diff --git a/sycl/unittests/helpers/PiImage.hpp b/sycl/unittests/helpers/PiImage.hpp index eee1fa4b20b6f..f55d4ccaaafe9 100644 --- a/sycl/unittests/helpers/PiImage.hpp +++ b/sycl/unittests/helpers/PiImage.hpp @@ -51,12 +51,13 @@ class PiProperty { private: void updateNativeType() { - if ( MType == PI_PROPERTY_TYPE_UINT32 ) { - MNative = NativeType{const_cast(MName.c_str()), - nullptr, MType, *((uint32_t*)MData.data())}; + if (MType == PI_PROPERTY_TYPE_UINT32) { + MNative = NativeType{const_cast(MName.c_str()), nullptr, MType, + *((uint32_t *)MData.data())}; } else { - MNative = NativeType{const_cast(MName.c_str()), - const_cast(MData.data()), MType, MData.size()}; + MNative = + NativeType{const_cast(MName.c_str()), + const_cast(MData.data()), MType, MData.size()}; } } std::string MName; @@ -374,7 +375,7 @@ void addSpecConstants(PiArray SpecConstants, /// Utility function to add ESIMD kernel flag to property set. void addESIMDFlag(PiPropertySet &Props) { std::vector ValData(sizeof(uint32_t)); - ValData[0]= 1; + ValData[0] = 1; PiProperty Prop{"isEsimdImage", ValData, PI_PROPERTY_TYPE_UINT32}; PiArray Value{std::move(Prop)}; diff --git a/sycl/unittests/misc/KernelBuildOptions.cpp b/sycl/unittests/misc/KernelBuildOptions.cpp index b763ca1e6efb1..f97c599deeaa0 100644 --- a/sycl/unittests/misc/KernelBuildOptions.cpp +++ b/sycl/unittests/misc/KernelBuildOptions.cpp @@ -42,7 +42,7 @@ static pi_result redefinedProgramCreate(pi_context, const void *, size_t, } static pi_result redefinedProgramBuild( - pi_program prog, pi_uint32, const pi_device *, const char * options, + pi_program prog, pi_uint32, const pi_device *, const char *options, void (*pfn_notify)(pi_program program, void *user_data), void *user_data) { if (options) BuildOpts = options; @@ -55,7 +55,7 @@ static pi_result redefinedProgramBuild( } static pi_result redefinedProgramCompile(pi_program, pi_uint32, - const pi_device *, const char * options, + const pi_device *, const char *options, pi_uint32, const pi_program *, const char **, void (*)(pi_program, void *), void *) { @@ -67,7 +67,7 @@ static pi_result redefinedProgramCompile(pi_program, pi_uint32, } static pi_result redefinedProgramLink(pi_context, pi_uint32, const pi_device *, - const char * options, pi_uint32, + const char *options, pi_uint32, const pi_program *, void (*)(pi_program, void *), void *, pi_program *) { @@ -168,9 +168,9 @@ static pi_result redefinedEnqueueKernelLaunch(pi_queue, pi_kernel, pi_uint32, static void setupDefaultMockAPIs(sycl::unittest::PiMock &Mock) { using namespace sycl::detail; Mock.redefine( - redefinedProgramCreateWithSource); + redefinedProgramCreateWithSource); Mock.redefine( - redefinedProgramCreateWithBinary); + redefinedProgramCreateWithBinary); Mock.redefine(redefinedProgramCreate); Mock.redefine(redefinedProgramCompile); Mock.redefine(redefinedProgramLink); @@ -243,7 +243,6 @@ TEST(KernelBuildOptions, KernelBundleBasic) { auto LinkBundle = sycl::link(ObjBundle, ObjBundle.get_devices()); // TODO: uncomment when image options are passed to BE // EXPECT_EQ(BuildOpts, "-link-img -vc-codegen"); - } TEST(KernelBuildOptions, Program) { @@ -270,8 +269,7 @@ TEST(KernelBuildOptions, Program) { sycl::program Prg2(Ctx); Prg1.build_with_source(""); - EXPECT_TRUE(BuildOpts.size()==0) << "Expect empty build options"; + EXPECT_TRUE(BuildOpts.size() == 0) << "Expect empty build options"; Prg2.build_with_source("", "-api-opts"); EXPECT_EQ(BuildOpts, "-api-opts"); - }