[Driver][SYCL] Provide ability to use an external host compiler

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -106,6 +106,8 @@ def err_drv_use_of_Z_option : Error<
   "unsupported use of internal gcc -Z option '%0'">;
 def err_drv_output_argument_with_multiple_files : Error<
   "cannot specify -o when generating multiple output files">;
+def err_drv_output_type_with_host_compiler : Error<
+  "unsupported output type when using external host compiler">;
 def err_drv_out_file_argument_with_multiple_sources : Error<
   "cannot specify '%0%1' when compiling multiple source files">;
 def err_no_external_assembler : Error<

clang/include/clang/Driver/Options.td

@@ -2518,6 +2518,13 @@ def fsycl_help : Flag<["-"], "fsycl-help">, Alias<fsycl_help_EQ>,
 def fsycl_libspirv_path_EQ : Joined<["-"], "fsycl-libspirv-path=">,
   Flags<[CC1Option, CoreOption]>, HelpText<"Path to libspirv library">;
 def fno_sycl_libspirv : Flag<["-"], "fno-sycl-libspirv">, HelpText<"Disable check for libspirv">;
+def fsycl_host_compiler_EQ : Joined<["-"], "fsycl-host-compiler=">,
+  Flags<[CoreOption]>, HelpText<"Specify C++ compiler binary to perform host "
+  "compilation with during SYCL offload compiles.">;
+def fsycl_host_compiler_options_EQ : Joined<["-"], "fsycl-host-compiler-options=">,
+  Flags<[CoreOption]>, HelpText<"When performing the host compilation with "
+  "-fsycl-host-compiler specified, use the given options during that compile. "
+  "Options are expected to be a quoted list of space separated options.">;
 def fsyntax_only : Flag<["-"], "fsyntax-only">,
   Flags<[NoXarchOption,CoreOption,CC1Option,FC1Option]>, Group<Action_Group>;
 def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group<f_Group>;

clang/lib/Driver/ToolChains/Clang.cpp

@@ -40,6 +40,7 @@
 #include "llvm/Option/ArgList.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/CodeGen.h"
+#include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/Support/Compression.h"
 #include "llvm/Support/FileSystem.h"
@@ -4192,6 +4193,167 @@ static bool ContainsWrapperAction(const Action *A) {
   return false;
 }
 
+// Put together an external compiler compilation call which is used instead
+// of the clang invocation for the host compile of an offload compilation.
+// Enabling command line:  clang++ -fsycl -fsycl-host-compiler=<HostExe>
+//                         <ClangOpts> -fsycl-host-compiler-options=<HostOpts>
+// Any <ClangOpts> used which are phase limiting (preprocessing, assembly,
+// object generation) are specifically handled here by specifying the
+// equivalent phase limiting option(s).
+// It is expected that any user <HostOpts> options passed will be placed
+// after any implied options set here.  This will have overriding behaviors
+// for any options which are considered to be evaluated from left to right.
+// Specifying any <HostOpts> option which conficts any of the implied options
+// will result in undefined behavior.  Potential conflicting options:
+//  * Output specification options (-o, -Fo, -Fa, etc)
+//  * Phase limiting options (-E, -c, -P, etc)
+void Clang::ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
+                                     const InputInfo &Output,
+                                     const InputInfoList &Inputs,
+                                     const llvm::opt::ArgList &TCArgs) const {
+
+  // The Host compilation step that occurs here is constructed based on the
+  // input from the user.  This consists of the compiler to call and the
+  // options that will be used during the compilation.
+  ArgStringList HostCompileArgs;
+  const InputInfo &InputFile = Inputs.front();
+  const ToolChain &TC = getToolChain();
+
+  // Input file.
+  HostCompileArgs.push_back(InputFile.getFilename());
+
+  // When performing the host compilation, we are expecting to only be
+  // creating intermediate files, namely preprocessor output, assembly or
+  // object files.
+  // We are making assumptions in regards to what options are used to
+  // generate these intermediate files.
+  //                gcc/g++/clang/clang++/default | cl
+  //  Object:                   -c                | -c
+  //  Preprocessed:             -E                | -P -Fi<file>
+  //  Assembly:                 -S                | -c -Fa<file>
+  //  Header Input:        -include <file>        | -FI <file>
+  //
+  // The options used are determined by the compiler name and target triple.
+  Arg *HostCompilerDefArg =
+      TCArgs.getLastArg(options::OPT_fsycl_host_compiler_EQ);
+  assert(HostCompilerDefArg && "Expected host compiler designation.");
+
+  bool OutputAdded = false;
+  StringRef CompilerName =
+      llvm::sys::path::stem(HostCompilerDefArg->getValue());
+  if (CompilerName.empty())
+    TC.getDriver().Diag(diag::err_drv_missing_arg_mtp)
+        << HostCompilerDefArg->getAsString(TCArgs);
+  // FIXME: Consider requiring user input to specify a compatibility class
+  // to determine the type of host compiler being used.
+  SmallVector<StringRef, 4> MSVCCompilers = {"cl", "clang-cl", "icl"};
+  bool IsMSVCHostCompiler =
+      std::find(MSVCCompilers.begin(), MSVCCompilers.end(), CompilerName) !=
+      MSVCCompilers.end();
+
+  auto addMSVCOutputFile = [&](StringRef Opt) {
+    SmallString<128> OutOpt(Opt);
+    OutOpt += Output.getFilename();
+    HostCompileArgs.push_back(TCArgs.MakeArgString(OutOpt));
+    OutputAdded = true;
+  };
+  // FIXME: Reuse existing toolchains which are already supported to put
+  // together the options.
+  // FIXME: For any potential obscure host compilers that do not use the
+  // 'standard' set of options, we should provide a user interface that allows
+  // users to override the implied options.
+  if (isa<PreprocessJobAction>(JA)) {
+    if (IsMSVCHostCompiler) {
+      // Check the output file, if it is 'stdout' we want to use -E.
+      if (StringRef(Output.getFilename()).equals("-")) {
+        HostCompileArgs.push_back("-E");
+        OutputAdded = true;
+      } else {
+        HostCompileArgs.push_back("-P");
+        addMSVCOutputFile("-Fi");
+      }
+    } else
+      HostCompileArgs.push_back("-E");
+  } else if (isa<AssembleJobAction>(JA)) {
+    HostCompileArgs.push_back("-c");
+    if (IsMSVCHostCompiler)
+      addMSVCOutputFile("-Fo");
+  } else {
+    assert((isa<CompileJobAction, BackendJobAction>(JA)) &&
+           "Invalid action for external host compilation tool.");
+    if (JA.getType() == types::TY_PP_Asm) {
+      if (IsMSVCHostCompiler) {
+        HostCompileArgs.push_back("-c");
+        addMSVCOutputFile("-Fa");
+        // The MSVC Compiler does not have a way to just create the assembly
+        // file so we create the assembly file and object file, and redirect
+        // the object file to a temporary.
+        std::string ObjTmpName = C.getDriver().GetTemporaryPath("host", "obj");
+        StringRef WrapperFileName =
+            C.addTempFile(C.getArgs().MakeArgString(ObjTmpName));
+        SmallString<128> ObjOutOpt("-Fo");
+        ObjOutOpt += WrapperFileName;
+        HostCompileArgs.push_back(C.getArgs().MakeArgString(ObjOutOpt));
+      } else
+        HostCompileArgs.push_back("-S");
+    } else {
+      TC.getDriver().Diag(diag::err_drv_output_type_with_host_compiler);
+    }
+  }
+
+  // Add default header search directories.
+  SmallString<128> BaseDir(C.getDriver().Dir);
+  llvm::sys::path::append(BaseDir, "..", "include");
+  SmallString<128> SYCLDir(BaseDir);
+  llvm::sys::path::append(SYCLDir, "sycl");
+  HostCompileArgs.push_back("-I");
+  HostCompileArgs.push_back(TCArgs.MakeArgString(SYCLDir));
+  HostCompileArgs.push_back("-I");
+  HostCompileArgs.push_back(TCArgs.MakeArgString(BaseDir));
+
+  if (!OutputAdded) {
+    // Add output file to the command line.  This is assumed to be prefaced
+    // with the '-o' option that is used to designate the output file.
+    HostCompileArgs.push_back("-o");
+    HostCompileArgs.push_back(Output.getFilename());
+  }
+
+  // Add the integration header.
+  StringRef Header =
+      TC.getDriver().getIntegrationHeader(InputFile.getBaseInput());
+  if (types::getPreprocessedType(InputFile.getType()) != types::TY_INVALID &&
+      !Header.empty()) {
+    HostCompileArgs.push_back(IsMSVCHostCompiler ? "-FI" : "-include");
+    HostCompileArgs.push_back(TCArgs.MakeArgString(Header));
+  }
+
+  SmallString<128> ExecPath;
+  if (HostCompilerDefArg) {
+    ExecPath = HostCompilerDefArg->getValue();
+    if (!ExecPath.empty() && ExecPath == llvm::sys::path::stem(ExecPath))
+      ExecPath = TC.GetProgramPath(ExecPath.c_str());
+  }
+
+  // Add any user-specified arguments.
+  if (Arg *HostCompilerOptsArg =
+          TCArgs.getLastArg(options::OPT_fsycl_host_compiler_options_EQ)) {
+    SmallVector<const char *, 8> TargetArgs;
+    llvm::BumpPtrAllocator BPA;
+    llvm::StringSaver S(BPA);
+    // Tokenize the string.
+    llvm::cl::TokenizeGNUCommandLine(HostCompilerOptsArg->getValue(), S,
+                                     TargetArgs);
+    llvm::transform(TargetArgs, std::back_inserter(HostCompileArgs),
+                    [&TCArgs](StringRef A) { return TCArgs.MakeArgString(A); });
+  }
+  const Tool *T = TC.SelectTool(JA);
+  auto Cmd = std::make_unique<Command>(JA, *T, ResponseFileSupport::None(),
+                                       TCArgs.MakeArgString(ExecPath),
+                                       HostCompileArgs, None);
+
+  C.addCommand(std::move(Cmd));
+}
+
 void Clang::ConstructJob(Compilation &C, const JobAction &JA,
                          const InputInfo &Output, const InputInfoList &Inputs,
                          const ArgList &Args, const char *LinkingOutput) const {
@@ -4225,6 +4387,14 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
           IsHeaderModulePrecompile || Inputs.size() == 1) &&
          "Unable to handle multiple inputs.");
 
+  // Perform the SYCL host compilation using an external compiler if the user
+  // requested.
+  if (Args.hasArg(options::OPT_fsycl_host_compiler_EQ) && IsSYCL &&
+      !IsSYCLOffloadDevice) {
+    ConstructHostCompilerJob(C, JA, Output, Inputs, Args);
+    return;
+  }
+
   // A header module compilation doesn't have a main input file, so invent a
   // fake one as a placeholder.
   const char *ModuleName = [&]{

clang/lib/Driver/ToolChains/Clang.h

@@ -88,6 +88,11 @@ class LLVM_LIBRARY_VISIBILITY Clang : public Tool {
                       codegenoptions::DebugInfoKind *DebugInfoKind,
                       bool *EmitCodeView) const;
 
+  void ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
+                                const InputInfo &Output,
+                                const InputInfoList &Inputs,
+                                const llvm::opt::ArgList &TCArgs) const;
+
   mutable std::unique_ptr<llvm::raw_fd_ostream> CompilationDatabase = nullptr;
   void DumpCompilationDatabase(Compilation &C, StringRef Filename,
                                StringRef Target,

clang/lib/Driver/ToolChains/MSVC.cpp

@@ -433,8 +433,10 @@ void visualstudio::Linker::ConstructJob(Compilation &C, const JobAction &JA,
     CmdArgs.push_back("-defaultlib:oldnames");
   }
 
-  if (!C.getDriver().IsCLMode() && !Args.hasArg(options::OPT_nostdlib) &&
-      Args.hasArg(options::OPT_fsycl) && !Args.hasArg(options::OPT_nolibsycl)) {
+  if ((!C.getDriver().IsCLMode() && !Args.hasArg(options::OPT_nostdlib) &&
+       Args.hasArg(options::OPT_fsycl) &&
+       !Args.hasArg(options::OPT_nolibsycl)) ||
+      Args.hasArg(options::OPT_fsycl_host_compiler_EQ)) {
     if (Args.hasArg(options::OPT__SLASH_MDd))
       CmdArgs.push_back("-defaultlib:sycld.lib");
     else

clang/test/Driver/sycl-host-compiler.cpp

@@ -0,0 +1,64 @@
+// Tests the abilities involved with using an external host compiler
+// REQUIRES: clang-driver
+
+/// enabling with -fsycl-host-compiler
+// RUN: %clangxx -fsycl -fsycl-host-compiler=/some/dir/g++ %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_COMPILER %s
+// HOST_COMPILER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]"
+// HOST_COMPILER: g++{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}} "-include" "[[INTHEADER]]"
+// HOST_COMPILER: ld{{.*}} "[[HOSTOBJ]]"
+
+// RUN: %clang_cl -fsycl -fsycl-host-compiler=/some/dir/cl %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_COMPILER_CL %s
+// HOST_COMPILER_CL: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]"
+// HOST_COMPILER_CL: cl{{.*}} "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-FI" "[[INTHEADER]]"
+// HOST_COMPILER_CL: link{{.*}} "[[HOSTOBJ]]"
+
+/// check for additional host options
+// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -fsycl-host-compiler-options="-DFOO -DBAR" %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_OPTIONS %s
+// HOST_OPTIONS: g++{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}} "-DFOO" "-DBAR"
+
+// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -fsycl-host-compiler-options="/DFOO /DBAR /O2" %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_OPTIONS_CL %s
+// HOST_OPTIONS_CL: cl{{.*}} "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "/DFOO" "/DBAR" "/O2"
+
+/// preprocessing
+// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -E %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_PREPROCESS %s
+// HOST_PREPROCESS: g++{{.*}} "-E"{{.*}} "-o" "[[PPOUT:.+\.ii]]"
+// HOST_PREPROCESS: g++{{.*}} "-E"{{.*}} "-o" "-"
+// HOST_PREPROCESS: clang-offload-bundler{{.*}} "-inputs={{.*}}.ii,[[PPOUT]]"
+
+// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -E %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_PREPROCESS_CL %s
+// HOST_PREPROCESS_CL: cl{{.*}} "-P"{{.*}} "-Fi[[PPOUT:.+\.ii]]"
+// HOST_PREPROCESS_CL: cl{{.*}} "-E"
+// HOST_PREPROCESS_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.ii,[[PPOUT]]"
+
+/// obj output
+// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -c %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_OBJECT %s
+// HOST_OBJECT: g++{{.*}} "-c"{{.*}} "-o" "[[OBJOUT:.+\.o]]"
+// HOST_OBJECT: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[OBJOUT]]"
+
+// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -c %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_OBJECT_CL %s
+// HOST_OBJECT_CL: cl{{.*}} "-c"{{.*}} "-Fo[[OBJOUT:.+\.obj]]"
+// HOST_OBJECT_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[OBJOUT]]"
+
+/// assembly output
+// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -S %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_ASSEMBLY %s
+// HOST_ASSEMBLY: g++{{.*}} "-S"{{.*}} "-o" "[[ASMOUT:.+\.s]]"
+// HOST_ASSEMBLY: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[ASMOUT]]"
+
+// RUN: %clangxx -fsycl -fsycl-host-compiler=cl -S %s -### 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_ASSEMBLY_CL %s
+// HOST_ASSEMBLY_CL: cl{{.*}} "-c"{{.*}} "-Fa[[ASMOUT:.+\.s]]" "-Fo{{.*}}.obj"
+// HOST_ASSEMBLY_CL: clang-offload-bundler{{.*}} "-inputs={{.*}}.bc,[[ASMOUT]]"
+
+/// missing argument error -fsycl-host-compiler=
+// RUN: %clangxx -fsycl -fsycl-host-compiler= -c -### %s 2>&1 \
+// RUN:  | FileCheck -check-prefix=HOST_COMPILER_NOARG %s
+// HOST_COMPILER_NOARG: missing argument to '-fsycl-host-compiler='

sycl/doc/UsersManual.md

@@ -222,6 +222,26 @@ and not recommended to use in production environment.
     same as specifying -fsycl-help with no argument and emits help for all
     backends.
 
+**`-fsycl-host-compiler=<arg>`**
+
+    Informs the compiler driver that the host compilation step that is performed
+    as part of the greater compilation flow will be performed by the compiler
+    <arg>.  It is expected that <arg> is the compiler to be called, either by
+    name (in which the PATH will be used to discover it) or a fully qualified
+    directory with compiler to invoke.  This option is only useful when -fsycl
+    is provided on the command line.
+
+**`-fsycl-host-compiler-options="opts"`**
+
+    Passes along the space separated quoted "opts" string as option arguments
+    to the compiler specified with the -fsycl-host-compiler=<arg> option.  It is
+    expected that the options used here are compatible with the compiler
+    specified via -fsycl-host-compiler=<arg>.
+
+    NOTE: Using -fsycl-host-compiler-options to pass any kind of phase limiting
+    options (e.g. -c, -E, -S) may interfere with the expected output set during
+    the host compilation.  Doing so is considered undefined behavior.
+
 # Example: SYCL device code compilation
 
 To invoke SYCL device compiler set `-fsycl-device-only` flag.

sycl/test/regression/fsycl-host-compiler-win.cpp

@@ -0,0 +1,46 @@
+// RUN: %clang_cl -fsycl -fsycl-host-compiler=cl -DDEFINE_CHECK -fsycl-host-compiler-options="-DDEFINE_CHECK" /Fe%t1.exe %s
+// RUN: %RUN_ON_HOST %t1.exe
+// REQUIRES: system-windows
+//
+//==------- fsycl-host-compiler-win.cpp - external host compiler 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Uses -fsycl-host-compiler=<compiler> on a simple test, requires 'cl'
+
+#include <CL/sycl.hpp>
+
+#ifndef DEFINE_CHECK
+#error predefined macro not set
+#endif // DEFINE_CHECK
+
+using namespace cl::sycl;
+
+int main() {
+  int data[] = {0, 0, 0};
+
+  {
+    buffer<int, 1> b(data, range<1>(3), {property::buffer::use_host_ptr()});
+    queue q;
+    q.submit([&](handler &cgh) {
+      auto B = b.get_access<access::mode::write>(cgh);
+      cgh.parallel_for<class test>(range<1>(3), [=](id<1> idx) {
+        B[idx] = 1;
+      });
+    });
+  }
+
+  bool isSuccess = true;
+
+  for (int i = 0; i < 3; i++)
+    if (data[i] != 1) isSuccess = false;
+
+  if (!isSuccess)
+    return -1;
+
+  return 0;
+}