[SYCL] Support -ftarget-compile-fast in JIT mode (#9516)

This option is currently only supported for AOT mode, and it's
implemented in the driver by passing an
`-igc_opts` string to ocloc.

Support this in JIT to by passing the FE option to
`clang-offload-wrapper` so it's stored in the binary and will get used
at runtime.

At runtime, check with the backend to get the correct backend flag, and
use that. If the device is not an Intel GPU, strip the FE flag.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -963,7 +963,6 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
   if (Arg *A = Args.getLastArg(options::OPT_O_Group))
     if (A->getOption().matches(options::OPT_O0))
       BeArgs.push_back("-cl-opt-disable");
-
   if (IsGen) {
     // For GEN (spir64_gen) we have implied -device settings given usage
     // of intel_gpu_ as a target.  Handle those here, and also check that no
@@ -991,10 +990,14 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
       CmdArgs.push_back("-device");
       CmdArgs.push_back(Args.MakeArgString(DepInfo));
     }
-    // -ftarget-compile-fast
+    // -ftarget-compile-fast AOT
     if (Args.hasArg(options::OPT_ftarget_compile_fast)) {
       BeArgs.push_back("-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'");
     }
+  } else if (Triple.getSubArch() == llvm::Triple::NoSubArch &&
+             Triple.isSPIR()) {
+    // -ftarget-compile-fast JIT
+    Args.AddLastArg(BeArgs, options::OPT_ftarget_compile_fast);
   }
   if (BeArgs.empty())
     return;

clang/test/Driver/ftarget-compile-fast.cpp

@@ -9,3 +9,12 @@
 
 // TARGET_COMPILE_FAST_GEN: ocloc{{.*}} "-output"
 // TARGET_COMPILE_FAST_GEN: "-options" "-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl \
+// RUN:    -ftarget-compile-fast %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=TARGET_COMPILE_FAST_JIT %s
+// RUN: %clang_cl -### --target=x86_64-pc-windows-msvc -fsycl \
+// RUN:     -ftarget-compile-fast %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=TARGET_COMPILE_FAST_JIT %s
+
+// TARGET_COMPILE_FAST_JIT: clang-offload-wrapper{{.*}} "-compile-opts=-ftarget-compile-fast

sycl/doc/UsersManual.md

@@ -410,6 +410,10 @@ and not recommended to use in production environment.
     Also, some of Intel GPUs or GPU run-time/drivers may support only
     "stateless" memory accesses.
 
+**`-ftarget-compile-fast`** [EXPERIMENTAL]
+    Instructs the target backend to reduce compilation time, potentially
+    at the cost of runtime performance. Currently only supported on Intel GPUs.
+
 # Example: SYCL device code compilation
 
 To invoke SYCL device compiler set `-fsycl-device-only` flag.

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -57,10 +57,11 @@ pi_result piPluginGetLastError(char **message) {
 }
 
 // Returns plugin specific backend option.
-// Current support is only for optimization options.
 // Return '-ze-opt-disable' for frontend_option = -O0.
 // Return '-ze-opt-level=1' for frontend_option = -O1 or -O2.
 // Return '-ze-opt-level=2' for frontend_option = -O3.
+// Return '-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'' for
+// frontend_option = -ftarget-compile-fast.
 pi_result piPluginGetBackendOption(pi_platform platform,
                                    const char *frontend_option,
                                    const char **backend_option) {

sycl/plugins/opencl/pi_opencl.cpp

@@ -99,8 +99,6 @@ pi_result piPluginGetLastError(char **message) {
 }
 
 // Returns plugin specific backend option.
-// Current support is only for optimization options.
-// Return '-cl-opt-disable' for frontend_option = -O0 and '' for others.
 pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option,
                                    const char **backend_option) {
   using namespace std::literals;
@@ -110,6 +108,7 @@ pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option,
     *backend_option = "";
     return PI_SUCCESS;
   }
+  // Return '-cl-opt-disable' for frontend_option = -O0 and '' for others.
   if (!strcmp(frontend_option, "-O0")) {
     *backend_option = "-cl-opt-disable";
     return PI_SUCCESS;
@@ -119,6 +118,10 @@ pi_result piPluginGetBackendOption(pi_platform, const char *frontend_option,
     *backend_option = "";
     return PI_SUCCESS;
   }
+  if (frontend_option == "-ftarget-compile-fast"sv) {
+    *backend_option = "-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'";
+    return PI_SUCCESS;
+  }
   return PI_ERROR_INVALID_VALUE;
 }
 

sycl/plugins/unified_runtime/ur/adapters/level_zero/ur_level_zero_platform.cpp

@@ -550,6 +550,8 @@ ur_result_t ur_platform_handle_t_::populateDeviceCacheIfNeeded() {
 // Return '-ze-opt-disable' for frontend_option = -O0.
 // Return '-ze-opt-level=1' for frontend_option = -O1 or -O2.
 // Return '-ze-opt-level=2' for frontend_option = -O3.
+// Return '-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'' for
+// frontend_option=-ftarget-compile-fast.
 UR_APIEXPORT ur_result_t UR_APICALL urPlatformGetBackendOption(
     ur_platform_handle_t Platform, ///< [in] handle of the platform instance.
     const char *FrontendOption, ///< [in] string containing the frontend option.
@@ -578,5 +580,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urPlatformGetBackendOption(
     *PlatformOption = "-ze-opt-level=2";
     return UR_RESULT_SUCCESS;
   }
+  if (FrontendOption == "-ftarget-compile-fast"sv) {
+    *PlatformOption = "-igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'";
+    return UR_RESULT_SUCCESS;
+  }
   return UR_RESULT_ERROR_INVALID_VALUE;
 }

sycl/source/detail/program_manager/program_manager.cpp

@@ -456,17 +456,34 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts,
       CompileOpts += std::string(backend_option);
     }
   }
-  if ((PlatformImpl->getBackend() == backend::ext_oneapi_level_zero ||
+  bool IsIntelGPU =
+      (PlatformImpl->getBackend() == backend::ext_oneapi_level_zero ||
        PlatformImpl->getBackend() == backend::opencl) &&
-      std::all_of(Devs.begin(), Devs.end(),
-                  [](const device &Dev) { return Dev.is_gpu(); }) &&
-      Img.getDeviceGlobals().size() != 0) {
+      std::all_of(Devs.begin(), Devs.end(), [](const device &Dev) {
+        return Dev.is_gpu() &&
+               Dev.get_info<info::device::vendor_id>() == 0x8086;
+      });
+  if (IsIntelGPU && Img.getDeviceGlobals().size() != 0) {
     // If the image has device globals we need to add the
     // -ze-take-global-address option to tell IGC to record addresses of these.
     if (!CompileOpts.empty())
       CompileOpts += " ";
     CompileOpts += "-ze-take-global-address";
   }
+  if (!CompileOptsEnv) {
+    static const char *TargetCompileFast = "-ftarget-compile-fast";
+    if (auto Pos = CompileOpts.find(TargetCompileFast);
+        Pos != std::string::npos) {
+      const char *BackendOption = nullptr;
+      if (IsIntelGPU)
+        PlatformImpl->getBackendOption(TargetCompileFast, &BackendOption);
+      auto OptLen = strlen(TargetCompileFast);
+      if (IsIntelGPU && BackendOption && BackendOption[0] != '\0')
+        CompileOpts.replace(Pos, OptLen, BackendOption);
+      else
+        CompileOpts.erase(Pos, OptLen);
+    }
+  }
 }
 
 static void applyOptionsFromImage(std::string &CompileOpts,

sycl/test-e2e/KernelAndProgram/target_compile_fast.cpp

@@ -0,0 +1,43 @@
+// RUN: %{build} -ftarget-compile-fast -o %t_with.out
+// RUN: %{build} -o %t_without.out
+
+// RUN: env SYCL_PI_TRACE=-1 %{run} %t_with.out 2>&1 | FileCheck %if !gpu || ext_oneapi_hip || ext_oneapi_cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-INTEL-GPU-WITH %} %s
+// RUN: env SYCL_PI_TRACE=-1 %{run} %t_without.out 2>&1 | FileCheck --implicit-check-not=-igc_opts %s
+
+// CHECK-INTEL-GPU-WITH: ---> piProgramBuild(
+// CHECK-INTEL-GPU-WITH: -igc_opts 'PartitionUnit=1,SubroutineThreshold=50000'
+
+// CHECK-WITHOUT: ---> piProgramBuild(
+// CHECK-WITHOUT-NOT: -igc_opts
+// CHECK-WITHOUT: ) ---> pi_result : PI_SUCCESS
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::buffer<size_t, 1> Buffer(4);
+
+  sycl::queue Queue;
+
+  sycl::range<1> NumOfWorkItems{Buffer.size()};
+
+  Queue.submit([&](sycl::handler &cgh) {
+    sycl::accessor Accessor{Buffer, cgh, sycl::write_only};
+    cgh.parallel_for<class FillBuffer>(NumOfWorkItems, [=](sycl::id<1> WIid) {
+      Accessor[WIid] = WIid.get(0);
+    });
+  });
+
+  sycl::host_accessor HostAccessor{Buffer, sycl::read_only};
+
+  bool MismatchFound = false;
+  for (size_t I = 0; I < Buffer.size(); ++I) {
+    if (HostAccessor[I] != I) {
+      std::cout << "The result is incorrect for element: " << I
+                << " , expected: " << I << " , got: " << HostAccessor[I]
+                << std::endl;
+      MismatchFound = true;
+    }
+  }
+
+  return MismatchFound;
+}