[SYCL] Add support for -ftarget-register-alloc-mode (#11430)

sarnex · web-flow · commit 8e5275af77af · 2023-10-16T21:09:06.000+02:00
We finally are using a new enough IGC driver in CI so we can implement
this. The option was approved a while back by the options group.

This change adds a new option `-ftarget-register-alloc-mode`. It is
specified as follows:

`-ftarget-register-alloc-mode=DeviceName0:Mode0[,DeviceName1:Mode1...]`

Currently the only valid device name is `pvc`, but I developed this
change to be extendable in the future.
The valid modes are:
auto -&gt; request the backend to use heuristics to pick a register alloc
mode
small -&gt; use small grf mode
large -&gt; use large grf mode
default -&gt; provide no specification to the backend on what register
alloc mode to use.

The default value if not specified is `pvc:default`, so provide no
specification to the backend in any case for any hardware. I will begin
an internal discussion after this is merged to see if we should change
the default to be `pvc:auto`.

The driver owns the mapping between the option mode and the backend
flag. It converts the user specified mode to the backend flag and passes
it either to ocloc for AOT or clang-offload-wrapper for JIT. The runtime
will extract the option stored in the image by clang-offload-wrapper and
pass it to the device backend if the device matches.

For AOT, we start making use of a new ocloc feature, `-device_options
&lt;device&gt; &lt;options&gt;`, ex `-device_options pvc "-options
-ze-opt-large-register-file"`.

This means "if the device is pvc, use large GRF mode. If the device is
not pvc, ignore the option."

For JIT, we do some argument processing in the runtime. The driver
splits the options in the case multiple are passed, and converts the
mode to the backend argument name to make the runtime processing easier.
For example, if the user passes
`-ftarget-target-alloc-mode=pvc:auto,pvc:large`, the options stored in
the image that the runtime sees will be
`-ftarget-target-alloc-mode=pvc:-ze-opt-large-register-file
-ftarget-target-alloc-mode=pvc:-ze-intel-enable-auto-large-GRF-mode`

The runtime loops over all the options, and for each instance of
`-ftarget-register-alloc-mode`:
Check if the device matches the device in the argument (currently only
pvc supported)
If so, replace the `-ftarget-register-alloc-mode=...` option with the
backend opt provided by the driver.
If not, remove the `-ftarget-register-alloc-mode=...` as to not pass it
to the backend.

---------

Signed-off-by: Sarnie, Nick &lt;nick.sarnie@intel.com&gt;
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
@@ -3840,6 +3840,10 @@ def ftarget_export_symbols : Flag<["-"], "ftarget-export-symbols">,
   "target library to allow for visibilty to other modules.">;
 def fno_target_export_symbols : Flag<["-"], "fno-target-export-symbols">,
   Visibility<[ClangOption, CLOption, DXCOption]>;
+def ftarget_register_alloc_mode_EQ : Joined<["-"], "ftarget-register-alloc-mode=">,
+  Visibility<[ClangOption, CLOption, DXCOption]>,
+  HelpText<"Specify a register allocation mode for specific hardware for use by supported "
+  "target backends.">;
 def : Flag<["-"], "fsycl-rdc">, Visibility<[ClangOption, CLOption, DXCOption]>, Alias<fgpu_rdc>;
 def : Flag<["-"], "fno-sycl-rdc">, Visibility<[ClangOption, CLOption, DXCOption]>, Alias<fno_gpu_rdc>;
 def fsycl_optimize_non_user_code : Flag<["-"], "fsycl-optimize-non-user-code">,
diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp
@@ -763,6 +763,14 @@ void SYCL::fpga::BackendCompiler::ConstructJob(
     C.addCommand(std::move(Cmd));
 }
 
+StringRef SYCL::gen::getGenGRFFlag(StringRef GRFMode) {
+  return llvm::StringSwitch<StringRef>(GRFMode)
+      .Case("auto", "-ze-intel-enable-auto-large-GRF-mode")
+      .Case("small", "-ze-intel-128-GRF-per-thread")
+      .Case("large", "-ze-opt-large-register-file")
+      .Default("");
+}
+
 void SYCL::gen::BackendCompiler::ConstructJob(Compilation &C,
                                               const JobAction &JA,
                                               const InputInfo &Output,
@@ -1132,6 +1140,9 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
   // GEN:  -options "-g -O0"
   // CPU:  "--bo=-g -cl-opt-disable"
   llvm::opt::ArgStringList BeArgs;
+  // Per-device argument vector storing the device name and the backend argument
+  // string
+  llvm::SmallVector<std::pair<StringRef, StringRef>, 16> PerDeviceArgs;
   bool IsGen = Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen;
   if (Arg *A = Args.getLastArg(options::OPT_g_Group, options::OPT__SLASH_Z7))
     if (!A->getOption().matches(options::OPT_g0))
@@ -1142,6 +1153,40 @@ 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");
+  StringRef RegAllocModeOptName = "-ftarget-register-alloc-mode=";
+  if (Arg *A = Args.getLastArg(options::OPT_ftarget_register_alloc_mode_EQ)) {
+    StringRef RegAllocModeVal = A->getValue(0);
+    auto ProcessElement = [&](StringRef Ele) {
+      auto [DeviceName, RegAllocMode] = Ele.split(':');
+      StringRef BackendOptName = SYCL::gen::getGenGRFFlag(RegAllocMode);
+      bool IsDefault = RegAllocMode.equals("default");
+      if (RegAllocMode.empty() || !DeviceName.equals("pvc") ||
+          (BackendOptName.empty() && !IsDefault)) {
+        getDriver().Diag(diag::err_drv_unsupported_option_argument)
+            << A->getSpelling() << Ele;
+      }
+      // "default" means "provide no specification to the backend", so
+      // we don't need to do anything here.
+      if (IsDefault)
+        return;
+      if (IsGen) {
+        // For AOT, Use ocloc's per-device options flag with the correct ocloc
+        // option to honor the user's specification.
+        PerDeviceArgs.push_back(
+            {DeviceName, Args.MakeArgString("-options " + BackendOptName)});
+      } else if (Triple.isSPIR() &&
+                 Triple.getSubArch() == llvm::Triple::NoSubArch) {
+        // For JIT, pass -ftarget-register-alloc-mode=Device:BackendOpt to
+        // clang-offload-wrapper to be processed by the runtime.
+        BeArgs.push_back(Args.MakeArgString(RegAllocModeOptName + DeviceName +
+                                            ":" + BackendOptName));
+      }
+    };
+    llvm::SmallVector<StringRef, 16> RegAllocModeArgs;
+    RegAllocModeVal.split(RegAllocModeArgs, ',');
+    for (StringRef Elem : RegAllocModeArgs)
+      ProcessElement(Elem);
+  }
   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
@@ -1181,6 +1226,13 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
     // -ftarget-compile-fast JIT
     Args.AddLastArg(BeArgs, options::OPT_ftarget_compile_fast);
   }
+  if (IsGen) {
+    for (auto [DeviceName, BackendArgStr] : PerDeviceArgs) {
+      CmdArgs.push_back("-device_options");
+      CmdArgs.push_back(Args.MakeArgString(DeviceName));
+      CmdArgs.push_back(Args.MakeArgString(BackendArgStr));
+    }
+  }
   if (BeArgs.empty())
     return;
   if (Triple.getSubArch() == llvm::Triple::NoSubArch ||
diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h
@@ -115,6 +115,7 @@ class LLVM_LIBRARY_VISIBILITY BackendCompiler : public Tool {
 
 StringRef resolveGenDevice(StringRef DeviceName);
 SmallString<64> getGenDeviceMacro(StringRef DeviceName);
+StringRef getGenGRFFlag(StringRef GRFMode);
 
 // // Prefix for GPU specific targets used for -fsycl-targets
 constexpr char IntelGPU[] = "intel_gpu_";
diff --git a/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp b/clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp
@@ -0,0 +1,92 @@
+// Test SYCL -ftarget-register-alloc-mode
+
+// RUN: %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:auto %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AUTO_AOT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:large %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=LARGE_AOT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:small %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=SMALL_AOT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:default %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=DEFAULT_AOT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:small,pvc:large %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=MULTIPLE_ARGS_AOT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:   -ftarget-register-alloc-mode=pvc:auto %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=AUTO_JIT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:   -ftarget-register-alloc-mode=pvc:large %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=LARGE_JIT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:   -ftarget-register-alloc-mode=pvc:small %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=SMALL_JIT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:   -ftarget-register-alloc-mode=pvc:default %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=DEFAULT_JIT %s
+
+// RUN: %clang -### -fsycl \
+// RUN:   -ftarget-register-alloc-mode=pvc:small,pvc:large %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=MULTIPLE_ARGS_JIT %s
+
+// RUN: not %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=dg2:auto %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=BAD_DEVICE %s
+
+// RUN: not %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=pvc:superlarge %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=BAD_MODE %s
+
+// RUN: not %clang -### -fsycl \
+// RUN:    -fsycl-targets=spir64_gen -ftarget-register-alloc-mode=dg2:superlarge %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=BAD_BOTH %s
+
+// AUTO_AOT: ocloc{{.*}} "-output"
+// AUTO_AOT: -device_options
+// AUTO_AOT: pvc
+// AUTO_AOT: "-options -ze-intel-enable-auto-large-GRF-mode"
+
+// LARGE_AOT: ocloc{{.*}} "-output"
+// LARGE_AOT: -device_options
+// LARGE_AOT: pvc
+// LARGE_AOT: "-options -ze-opt-large-register-file"
+
+// SMALL_AOT: ocloc{{.*}} "-output"
+// SMALL_AOT: -device_options
+// SMALL_AOT: pvc
+// SMALL_AOT: "-options -ze-intel-128-GRF-per-thread"
+
+// DEFAULT_AOT-NOT: -device_options
+
+// MULTIPLE_ARGS_AOT: ocloc{{.*}} "-output"
+// MULTIPLE_ARGS_AOT: -device_options
+// MULTIPLE_ARGS_AOT: pvc
+// MULTIPLE_ARGS_AOT: "-options -ze-intel-128-GRF-per-thread"
+// MULTIPLE_ARGS_AOT: -device_options
+// MULTIPLE_ARGS_AOT: pvc
+// MULTIPLE_ARGS_AOT: "-options -ze-opt-large-register-file"
+
+// AUTO_JIT: clang-offload-wrapper{{.*}} "-compile-opts=-ftarget-register-alloc-mode=pvc:-ze-intel-enable-auto-large-GRF-mode"
+
+// LARGE_JIT: clang-offload-wrapper{{.*}} "-compile-opts=-ftarget-register-alloc-mode=pvc:-ze-opt-large-register-file"
+
+// SMALL_JIT: clang-offload-wrapper{{.*}} "-compile-opts=-ftarget-register-alloc-mode=pvc:-ze-intel-128-GRF-per-thread"
+
+// DEFAULT_JIT-NOT: -ftarget-register-alloc-mode=
+
+// MULTIPLE_ARGS_JIT: clang-offload-wrapper{{.*}} "-compile-opts=-ftarget-register-alloc-mode=pvc:-ze-intel-128-GRF-per-thread -ftarget-register-alloc-mode=pvc:-ze-opt-large-register-file"
+
+// BAD_DEVICE: unsupported argument 'dg2:auto' to option '-ftarget-register-alloc-mode='
+// BAD_MODE: unsupported argument 'pvc:superlarge' to option '-ftarget-register-alloc-mode='
+// BAD_BOTH: unsupported argument 'dg2:superlarge' to option '-ftarget-register-alloc-mode='
diff --git a/sycl/doc/UsersManual.md b/sycl/doc/UsersManual.md
@@ -426,6 +426,7 @@ and not recommended to use in production environment.
     "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.
 
@@ -436,6 +437,13 @@ and not recommended to use in production environment.
 
     NOTE: This flag is only supported for spir64_gen AOT targets.
 
+**`-ftarget-register-alloc-mode=<arg>`**
+
+    Specify a register allocation mode for specific hardware for use by supported
+    target backends. The format of the argument is "Device0:Mode0[,Device1:Mode1...]".
+    Currently the only supported Device is "pvc". The supported modes are
+    "default","small","large", and "auto".
+
 # Example: SYCL device code compilation
 
 To invoke SYCL device compiler set `-fsycl-device-only` flag.
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
@@ -494,6 +494,37 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts,
       else
         CompileOpts.erase(Pos, OptLen);
     }
+    static const std::string TargetRegisterAllocMode =
+        "-ftarget-register-alloc-mode=";
+    auto OptPos = CompileOpts.find(TargetRegisterAllocMode);
+    while (OptPos != std::string::npos) {
+      auto EndOfOpt = CompileOpts.find(" ", OptPos);
+      // Extract everything after the equals until the end of the option
+      auto OptValue = CompileOpts.substr(
+          OptPos + TargetRegisterAllocMode.size(),
+          EndOfOpt - OptPos - TargetRegisterAllocMode.size());
+      auto ColonPos = OptValue.find(":");
+      auto Device = OptValue.substr(0, ColonPos);
+      std::string BackendStrToAdd;
+      bool IsPVC =
+          std::all_of(Devs.begin(), Devs.end(), [&](const device &Dev) {
+            return IsIntelGPU &&
+                   (Dev.get_info<ext::intel::info::device::device_id>() &
+                    0xFF00) == 0x0B00;
+          });
+      // Currently 'pvc' is the only supported device.
+      if (Device == "pvc" && IsPVC)
+        BackendStrToAdd = " " + OptValue.substr(ColonPos + 1) + " ";
+
+      // Extract everything before this option
+      std::string NewCompileOpts =
+          CompileOpts.substr(0, OptPos) + BackendStrToAdd;
+      // Extract everything after this option and add it to the above.
+      if (EndOfOpt != std::string::npos)
+        NewCompileOpts += CompileOpts.substr(EndOfOpt);
+      CompileOpts = NewCompileOpts;
+      OptPos = CompileOpts.find(TargetRegisterAllocMode);
+    }
   }
 }
 
diff --git a/sycl/test-e2e/KernelAndProgram/target_register_alloc_mode.cpp b/sycl/test-e2e/KernelAndProgram/target_register_alloc_mode.cpp
@@ -0,0 +1,41 @@
+// REQUIRES: gpu-intel-pvc
+
+// RUN: %{build} -ftarget-register-alloc-mode=pvc:auto -o %t_with.out
+// RUN: %{build} -o %t_without.out
+
+// RUN: env SYCL_PI_TRACE=-1 %{run} %t_with.out 2>&1 | FileCheck --check-prefix=CHECK-WITH %s
+// RUN: env SYCL_PI_TRACE=-1 %{run} %t_without.out 2>&1 | FileCheck --implicit-check-not=-ze-intel-enable-auto-large-GRF-mode %s
+
+// CHECK-WITH: ---> piProgramBuild(
+// CHECK-WITH: -ze-intel-enable-auto-large-GRF-mode
+
+#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;
+}
