From a0117ab8692a67598046620e8a61adca3e6eb9a8 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Tue, 28 Nov 2023 19:44:10 +0000 Subject: [PATCH] [SYCL] Add internal debugging environment variables to append to compile/link options (#11981) Add two new environment variables to append to compile or link options. This is only intended for SYCL developers, and end users should use the '-X' family of dpcpp options. We are adding this because we have had some people hit unexpected issues because if `SYCL_PROGRAM_COMPILE_OPTIONS` is set, any options that would normally be added internally by the runtime and invisibly to the user (`ze-opt-disable`, `ze-opt-level`, etc) are not added, so someone doing A/B testing would not be comparing apples to apples with the envvar set vs not set. --------- Signed-off-by: Sarnie, Nick --- sycl/doc/EnvironmentVariables.md | 6 ++-- sycl/source/detail/config.def | 2 ++ .../program_manager/program_manager.cpp | 35 +++++++++++++++++-- sycl/test-e2e/Config/env_vars.cpp | 2 ++ .../kernel-bundle-merge-options-env.cpp | 10 +++--- 5 files changed, 46 insertions(+), 9 deletions(-) diff --git a/sycl/doc/EnvironmentVariables.md b/sycl/doc/EnvironmentVariables.md index 4918c554460ad..2a94326756c19 100644 --- a/sycl/doc/EnvironmentVariables.md +++ b/sycl/doc/EnvironmentVariables.md @@ -227,8 +227,10 @@ variables in production code. | `SYCL_DISABLE_EXECUTION_GRAPH_CLEANUP` | Any(\*) | Disable regular cleanup of enqueued (or finished, in case of host tasks) non-leaf command nodes. If disabled, command nodes will be cleaned up only during the destruction of the last remaining memory object used by them. | | `SYCL_DISABLE_POST_ENQUEUE_CLEANUP` (deprecated) | Any(\*) | Use `SYCL_DISABLE_EXECUTION_GRAPH_CLEANUP` instead. | | `SYCL_DEVICELIB_INHIBIT_NATIVE` | String of device library extensions (separated by a whitespace) | Do not rely on device native support for devicelib extensions listed in this option. | -| `SYCL_PROGRAM_COMPILE_OPTIONS` | String of valid OpenCL compile options | Override compile options for all programs. | -| `SYCL_PROGRAM_LINK_OPTIONS` | String of valid OpenCL link options | Override link options for all programs. | +| `SYCL_PROGRAM_COMPILE_OPTIONS` | String of valid compile options | Override compile options for all programs. | +| `SYCL_PROGRAM_LINK_OPTIONS` | String of valid link options | Override link options for all programs. | +| `SYCL_PROGRAM_APPEND_COMPILE_OPTIONS` | String of valid compile options | Append to the end of compile options for all programs. | +| `SYCL_PROGRAM_APPEND_LINK_OPTIONS` | String of valid link options | Append to the end of link options for all programs. | | `SYCL_USE_KERNEL_SPV` | Path to the SPIR-V binary | Load device image from the specified file. If runtime is unable to read the file, `sycl::runtime_error` exception is thrown. The image is assumed to have been created using the `-fno-sycl-dead-args-optimization` option. | | `SYCL_DUMP_IMAGES` | Any(\*) | Dump device image binaries to file. Control has no effect if `SYCL_USE_KERNEL_SPV` is set. | | `SYCL_HOST_UNIFIED_MEMORY` | Integer | Enforce host unified memory support or lack of it for the execution graph builder. If set to 0, it is enforced as not supported by all devices. If set to 1, it is enforced as supported by all devices. | diff --git a/sycl/source/detail/config.def b/sycl/source/detail/config.def index 69a1df48dc0c4..33d43934c66eb 100644 --- a/sycl/source/detail/config.def +++ b/sycl/source/detail/config.def @@ -22,6 +22,8 @@ CONFIG(SYCL_DEVICELIB_NO_FALLBACK, 1, __SYCL_DEVICELIB_NO_FALLBACK) CONFIG(SYCL_DEVICE_FILTER, 1024, __SYCL_DEVICE_FILTER) CONFIG(SYCL_PROGRAM_LINK_OPTIONS, 64, __SYCL_PROGRAM_LINK_OPTIONS) CONFIG(SYCL_PROGRAM_COMPILE_OPTIONS, 64, __SYCL_PROGRAM_COMPILE_OPTIONS) +CONFIG(SYCL_PROGRAM_APPEND_LINK_OPTIONS, 64, __SYCL_PROGRAM_APPEND_LINK_OPTIONS) +CONFIG(SYCL_PROGRAM_APPEND_COMPILE_OPTIONS, 64, __SYCL_PROGRAM_APPEND_COMPILE_OPTIONS) CONFIG(SYCL_HOST_UNIFIED_MEMORY, 1, __SYCL_HOST_UNIFIED_MEMORY) // 260 (Windows limit) - 12 (filename) - 84 (cache directory structure) CONFIG(SYCL_CACHE_DIR, 164, __SYCL_CACHE_DIR) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1bbfeabc449a0..2ebeaa05b6023 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -530,6 +530,26 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts, } } +static void +appendCompileEnvironmentVariablesThatAppend(std::string &CompileOpts) { + static const char *AppendCompileOptsEnv = + SYCLConfig::get(); + if (AppendCompileOptsEnv) { + if (!CompileOpts.empty()) + CompileOpts += " "; + CompileOpts += AppendCompileOptsEnv; + } +} +static void appendLinkEnvironmentVariablesThatAppend(std::string &LinkOpts) { + static const char *AppendLinkOptsEnv = + SYCLConfig::get(); + if (AppendLinkOptsEnv) { + if (!LinkOpts.empty()) + LinkOpts += " "; + LinkOpts += AppendLinkOptsEnv; + } +} + static void applyOptionsFromImage(std::string &CompileOpts, std::string &LinkOpts, const RTDeviceBinaryImage &Img, @@ -646,7 +666,9 @@ sycl::detail::pi::PiProgram ProgramManager::getBuiltPIProgram( &LinkOpts, SpecConsts] { const PluginPtr &Plugin = ContextImpl->getPlugin(); applyOptionsFromImage(CompileOpts, LinkOpts, Img, {Device}, Plugin); - + // Should always come last! + appendCompileEnvironmentVariablesThatAppend(CompileOpts); + appendLinkEnvironmentVariablesThatAppend(LinkOpts); auto [NativePrg, DeviceCodeWasInCache] = getOrCreatePIProgram( Img, Context, Device, CompileOpts + LinkOpts, SpecConsts); @@ -734,6 +756,9 @@ ProgramManager::getOrCreateKernel(const ContextImplPtr &ContextImpl, std::string CompileOpts, LinkOpts; SerializedObj SpecConsts; applyOptionsFromEnvironment(CompileOpts, LinkOpts); + // Should always come last! + appendCompileEnvironmentVariablesThatAppend(CompileOpts); + appendLinkEnvironmentVariablesThatAppend(LinkOpts); const sycl::detail::pi::PiDevice PiDevice = DeviceImpl->getHandleRef(); auto key = std::make_tuple(std::move(SpecConsts), PiDevice, @@ -2131,6 +2156,8 @@ ProgramManager::compile(const device_image_plain &DeviceImage, applyCompileOptionsFromEnvironment(CompileOptions); appendCompileOptionsFromImage( CompileOptions, *(InputImpl->get_bin_image_ref()), Devs, Plugin); + // Should always come last! + appendCompileEnvironmentVariablesThatAppend(CompileOptions); sycl::detail::pi::PiResult Error = Plugin->call_nocheck( ObjectImpl->get_program_ref(), /*num devices=*/Devs.size(), @@ -2169,6 +2196,8 @@ ProgramManager::link(const device_image_plain &DeviceImage, appendLinkOptionsFromImage(LinkOptionsStr, *(InputImpl->get_bin_image_ref())); } + // Should always come last! + appendLinkEnvironmentVariablesThatAppend(LinkOptionsStr); const context &Context = getSyclObjImpl(DeviceImage)->get_context(); const ContextImplPtr ContextImpl = getSyclObjImpl(Context); const PluginPtr &Plugin = ContextImpl->getPlugin(); @@ -2279,7 +2308,9 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage, ContextImplPtr ContextImpl = getSyclObjImpl(Context); const PluginPtr &Plugin = ContextImpl->getPlugin(); applyOptionsFromImage(CompileOpts, LinkOpts, Img, Devs, Plugin); - + // Should always come last! + appendCompileEnvironmentVariablesThatAppend(CompileOpts); + appendLinkEnvironmentVariablesThatAppend(LinkOpts); // TODO: Add support for creating non-SPIRV programs from multiple devices. if (InputImpl->get_bin_image_ref()->getFormat() != PI_DEVICE_BINARY_TYPE_SPIRV && diff --git a/sycl/test-e2e/Config/env_vars.cpp b/sycl/test-e2e/Config/env_vars.cpp index ff144df6101ea..2cf3f2efc2fcb 100644 --- a/sycl/test-e2e/Config/env_vars.cpp +++ b/sycl/test-e2e/Config/env_vars.cpp @@ -4,12 +4,14 @@ // RUN: %{build} -O0 -o %t.out // // RUN: env SYCL_PROGRAM_COMPILE_OPTIONS="-g" %{run} %t.out +// RUN: env SYCL_PROGRAM_APPEND_COMPILE_OPTIONS="-g" %{run} %t.out // // Now test for invalid options to make sure they are really passed to // a device compiler. Intel GPU runtime doesn't give an error for // invalid options, so we don't test it here. // // RUN: %if cpu %{ env SYCL_PROGRAM_COMPILE_OPTIONS="-enable-link-options -cl-denorms-are-zero" SHOULD_CRASH=1 %{run} %t.out %} +// RUN: %if cpu %{ env SYCL_PROGRAM_APPEND_COMPILE_OPTIONS="-enable-link-options -cl-denorms-are-zero" SHOULD_CRASH=1 %{run} %t.out %} #include #include diff --git a/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options-env.cpp b/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options-env.cpp index bf15b63ab2f73..37edd8f47ebfd 100644 --- a/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options-env.cpp +++ b/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options-env.cpp @@ -1,10 +1,10 @@ // REQUIRES: gpu // Disable fallback assert here so, that build process isn't affected // RUN: %{build} -DSYCL_DISABLE_FALLBACK_ASSERT=1 -o %t.out %debug_option -// RUN: env SYCL_PI_TRACE=-1 SYCL_PROGRAM_COMPILE_OPTIONS=-DENV_COMPILE_OPTS SYCL_PROGRAM_LINK_OPTIONS=-DENV_LINK_OPTS %{run} %t.out | FileCheck %s +// RUN: env SYCL_PI_TRACE=-1 SYCL_PROGRAM_COMPILE_OPTIONS=-DENV_COMPILE_OPTS SYCL_PROGRAM_LINK_OPTIONS=-DENV_LINK_OPTS SYCL_PROGRAM_APPEND_COMPILE_OPTIONS=-DENV_APPEND_COMPILE_OPTS SYCL_PROGRAM_APPEND_LINK_OPTIONS=-DENV_APPEND_LINK_OPTS %{run} %t.out | FileCheck %s // Check that options are overrided // RUN: %{build} -DSYCL_DISABLE_FALLBACK_ASSERT=1 -Xsycl-target-linker=spir64 -DBAR -Xsycl-target-frontend=spir64 -DBAR_COMPILE -o %t.out -// RUN: env SYCL_PI_TRACE=-1 SYCL_PROGRAM_COMPILE_OPTIONS=-DENV_COMPILE_OPTS SYCL_PROGRAM_LINK_OPTIONS=-DENV_LINK_OPTS %{run} %t.out | FileCheck %s +// RUN: env SYCL_PI_TRACE=-1 SYCL_PROGRAM_COMPILE_OPTIONS=-DENV_COMPILE_OPTS SYCL_PROGRAM_LINK_OPTIONS=-DENV_LINK_OPTS SYCL_PROGRAM_APPEND_COMPILE_OPTIONS=-DENV_APPEND_COMPILE_OPTS SYCL_PROGRAM_APPEND_LINK_OPTIONS=-DENV_APPEND_LINK_OPTS %{run} %t.out | FileCheck %s // UNSUPPORTED: hip #include "kernel-bundle-merge-options.hpp" @@ -13,16 +13,16 @@ // CHECK-NEXT: // CHECK-NEXT: // CHECK-NEXT: -// CHECK: :{{[^bar]*}}-DENV_COMPILE_OPTS{{[^bar]*}}-DENV_LINK_OPTS{{[^bar]*}} +// CHECK: :{{[^bar]*}}-DENV_COMPILE_OPTS -DENV_APPEND_COMPILE_OPTS{{[^bar]*}}-DENV_LINK_OPTS -DENV_APPEND_LINK_OPTS{{[^bar]*}} // CHECK: piProgramCompile( // CHECK-NEXT: // CHECK-NEXT: // CHECK-NEXT: -// CHECK: :{{[^bar]*}}-DENV_COMPILE_OPTS{{[^bar]*}} +// CHECK: :{{[^bar]*}}-DENV_COMPILE_OPTS -DENV_APPEND_COMPILE_OPTS{{[^bar]*}} // CHECK: piProgramLink( // CHECK-NEXT: // CHECK-NEXT: // CHECK-NEXT: -// CHECK: :{{[^bar]*}}-DENV_LINK_OPTS{{[^bar]*}} +// CHECK: :{{[^bar]*}}-DENV_LINK_OPTS -DENV_APPEND_LINK_OPTS{{[^bar]*}}