[SYCL][NewOffloadModel] Add ESIMD Procesing to clang-linker-wrapper.

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -728,6 +728,32 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args,
     PostLinkArgs.push_back("-generate-device-image-default-spec-consts");
 }
 
+static sycl::ESIMDProcessingOptions
+getTripleBasedESIMDPostLinkOpts(const ArgList &Args,
+                                const llvm::Triple Triple) {
+  sycl::ESIMDProcessingOptions Options;
+
+  // On Intel targets we don't need non-kernel functions as entry points,
+  // because it only increases amount of code for device compiler to handle,
+  // without any actual benefits.
+  // TODO: Try to extend this feature for non-Intel GPUs.
+  if (!Args.hasFlag(OPT_no_sycl_remove_unused_external_funcs,
+                    OPT_sycl_remove_unused_external_funcs, false) &&
+      !Args.hasArg(OPT_sycl_allow_device_image_dependencies) &&
+      !Triple.isNVPTX() && !Triple.isAMDGPU())
+    Options.EmitOnlyKernelsAsEntryPoints = true;
+
+  bool SplitEsimdByDefault = Triple.isSPIROrSPIRV();
+  bool SplitEsimd =
+      Args.hasFlag(OPT_sycl_device_code_split_esimd,
+                   OPT_no_sycl_device_code_split_esimd, SplitEsimdByDefault);
+  if (SplitEsimd)
+    Options.SplitESIMD = true;
+
+  Options.LowerESIMD = true;
+  return Options;
+}
+
 /// Run sycl-post-link tool for SYCL offloading.
 /// 'InputFiles' is the list of input LLVM IR files.
 /// 'Args' encompasses all arguments required for linking and wrapping device
@@ -832,6 +858,8 @@ runSYCLPostLinkLibrary(ArrayRef<StringRef> InputFiles, const ArgList &Args,
       IsAOTTarget)
     Settings.GenerateModuleDescWithDefaultSpecConsts = true;
 
+  Settings.ESIMDOptions = getTripleBasedESIMDPostLinkOpts(Args, Triple);
+
   if (DryRun) {
     auto OutputFileOrErr = createOutputFile(
         sys::path::filename(ExecutableName) + ".sycl.split.image", "bc");

llvm/include/llvm/SYCLPostLink/ESIMDPostSplitProcessing.h

@@ -33,6 +33,9 @@ struct ESIMDProcessingOptions {
   bool ForceDisableESIMDOpt = false;
 };
 
+SmallString<0>
+convertESIMDOptionsToString(const ESIMDProcessingOptions &Options);
+
 /// Lowers ESIMD constructs after separation from regular SYCL code.
 /// \p Options.SplitESIMD identifies that ESIMD splitting is requested in the
 /// compilation. Returns true if the given \p MD has been modified.

llvm/include/llvm/SYCLPostLink/SYCLPostLink.h

@@ -70,6 +70,8 @@ struct PostLinkSettings {
 
   std::optional<SpecConstantsPass::HandlingMode> SpecConstMode;
   bool GenerateModuleDescWithDefaultSpecConsts = false;
+
+  llvm::sycl::ESIMDProcessingOptions ESIMDOptions;
 };
 
 std::string convertSettingsToString(const PostLinkSettings &Settings);

llvm/lib/SYCLPostLink/ESIMDPostSplitProcessing.cpp

@@ -84,6 +84,20 @@ linkModules(std::unique_ptr<ModuleDesc> MD1, std::unique_ptr<ModuleDesc> MD2) {
 
 } // anonymous namespace
 
+SmallString<0>
+sycl::convertESIMDOptionsToString(cosnt sycl::ESIMDProcessingOptions &Options) {
+  return formatv(
+             "esimd.split_mode: {0}, esimd.EmitOnlyKernelsAsEntryPoints: {1}, "
+             "esimd.AllowDeviceImageDependencies: {2}, esimd.LowerESIMD: {3}, "
+             "esimd.SplitESIMD: {4}, esimd.OptLevel: {5}, "
+             "esimd.ForceDisableESIMDOpt: {6}",
+             module_split::convertSplitModeToString(Options.SplitMode),
+             Options.EmitOnlyKernelsAsEntryPoints,
+             Options.AllowDeviceImageDependencies, Options.LowerESIMD,
+             Options.SplitESIMD, Options.OptLevel, Options.ForceDisableESIMDOpt)
+      .sstr<0>();
+}
+
 // When ESIMD code was separated from the regular SYCL code,
 // we can safely process ESIMD part.
 bool sycl::lowerESIMDConstructs(ModuleDesc &MD,

llvm/lib/SYCLPostLink/SYCLPostLink.cpp

@@ -15,6 +15,7 @@
 #include "llvm/IR/Module.h"
 #include "llvm/IR/PassInstrumentation.h"
 #include "llvm/IRPrinter/IRPrintingPasses.h"
+#include "llvm/SYCLPostLink/ESIMDPostSplitProcessing.h"
 #include "llvm/Support/Error.h"
 #include "llvm/Support/FileSystem.h"
 #include "llvm/Support/FormatVariadic.h"
@@ -124,10 +125,12 @@ std::string llvm::sycl_post_link::convertSettingsToString(
 
   return formatv(
       "output_assembly: {0}, split_mode: {1}, specialization_constant_mode: "
-      "{2}, generate_module_with_default_spec_const_values: {3}",
+      "{2}, generate_module_with_default_spec_const_values: {3}, "
+      "{4}",
       Settings.OutputAssembly,
       module_split::convertSplitModeToString(Settings.SplitMode), SpecConstMode,
-      Settings.GenerateModuleDescWithDefaultSpecConsts);
+      Settings.GenerateModuleDescWithDefaultSpecConsts,
+      sycl::convertESIMDOptionsToString(Settings.ESIMDOptions));
 }
 
 Expected<std::vector<module_split::SplitModule>>
@@ -139,10 +142,19 @@ llvm::sycl_post_link::performPostLinkProcessing(
       [&SplitModules,
        Settings](std::unique_ptr<module_split::ModuleDesc> M) -> Error {
     M->fixupLinkageOfDirectInvokeSimdTargets();
-    // TODO: add ESIMD handling.
 
-    SmallVector<std::unique_ptr<module_split::ModuleDesc>> Modules;
-    Modules.push_back(std::move(M));
+    bool Modified = false;
+    bool SplitOccurred = false;
+    auto ModulesOrErr = sycl::handleESIMD(std::move(M), Settings.ESIMDOptions,
+                                          Modified, SplitOccurred);
+    if (!ModulesOrErr)
+      return ModulesOrErr.takeError();
+
+    SmallVector<std::unique_ptr<module_split::ModuleDesc>, 2> &Modules =
+        *ModulesOrErr;
+    assert(Modules.size() &&
+           "at least one module is expected after ESIMD split");
+
     SmallVector<std::unique_ptr<module_split::ModuleDesc>> NewModules;
     if (Settings.SpecConstMode)
       llvm::sycl_post_link::handleSpecializationConstants(
@@ -154,8 +166,9 @@ llvm::sycl_post_link::performPostLinkProcessing(
 
     for (std::unique_ptr<module_split::ModuleDesc> &MD : Modules) {
       size_t ID = SplitModules.size();
+      StringRef Suffix = MD.isESIMD() ? "_esimd" : "";
       std::string OutIRFilename =
-          (Settings.OutputPrefix + "_" + Twine(ID)).str();
+          (Settings.OutputPrefix + Suffix + "_" + Twine(ID)).str();
       Expected<module_split::SplitModule> SplitImageOrErr =
           saveModuleDesc(*MD, OutIRFilename, Settings.OutputAssembly);
       if (!SplitImageOrErr)