[SYCL][NewOffload][NFC] Move post-link processing to library for later reuse in NewOffload Model. (#17357)

This patch extracts post-link/pre-split processing into the dedicated
library function for later reuse in
clang-sycl-linker/clang-linker-wrapper.

The functions `removeSYCLKernelsConstRefArray` and
`removeDeviceGlobalFromCompilerUsed` have been transformed into LLVM
Passes.

Passes that were previously run individually are now executed together
by the PassManager. This change allows to introduce the use of
PassPipeline's debugging capabilities (e.g., -print-passes) in
NewOffload Model later.

In cases where the `invoke_simd` builtin is used in conjunction with
`-split-esimd`, a straightforward check has been introduced instead of
the previous post-processing check in sycl-post-link. This scenario
should also be checked in the NewOffload Model.

Author: maksimsab

llvm/include/llvm/SYCLLowerIR/CleanupSYCLMetadata.h

@@ -23,6 +23,33 @@ class CleanupSYCLMetadataPass : public PassInfoMixin<CleanupSYCLMetadataPass> {
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
 };
 
+/// Removes the global variable "llvm.used".
+/// "llvm.used" is a global constant array containing references to kernels
+/// available in the module and callable from host code. The elements of
+/// the array are ConstantExpr bitcast to i8*.
+/// The variable must be removed because it has done the job to the moment
+/// of a compilation stage and the references to the kernels callable from
+/// host must not have users.
+class CleanupSYCLMetadataFromLLVMUsed
+    : public PassInfoMixin<CleanupSYCLMetadataFromLLVMUsed> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
+/// Removes all device_global variables from the llvm.compiler.used global
+/// variable. A device_global with internal linkage will be in
+/// llvm.compiler.used to avoid the compiler wrongfully removing it during
+/// optimizations. However, as an effect the device_global variables will also
+/// be distributed across binaries, even if llvm.compiler.used has served its
+/// purpose. To avoid polluting other binaries with unused device_global
+/// variables, we remove them from llvm.compiler.used and erase them if they
+/// have no further uses.
+class RemoveDeviceGlobalFromLLVMCompilerUsed
+    : public PassInfoMixin<RemoveDeviceGlobalFromLLVMCompilerUsed> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
 } // namespace llvm
 
 #endif // LLVM_CLEANUP_SYCL_METADATA

llvm/include/llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h

@@ -46,6 +46,8 @@ bool isESIMDKernel(const Function &F);
 bool isESIMD(const Function &F);
 // Tells whether given function is a kernel.
 bool isKernel(const Function &F);
+// Tells whether a given Module contains an invoke_simd builtin.
+bool moduleContainsInvokeSimdBuiltin(Module &M);
 
 /// Reports and error with the message \p Msg concatenated with the optional
 /// \p OptMsg if \p Condition is false.

llvm/include/llvm/SYCLLowerIR/ModuleSplitter.h

@@ -333,6 +333,11 @@ struct ModuleSplitterSettings {
 /// Parses the output table file from sycl-post-link tool.
 Expected<std::vector<SplitModule>> parseSplitModulesFromFile(StringRef File);
 
+/// PreSplitProcessingPipeline maintains correctness.
+/// Note: After migration from sycl-post-link to NewOffload Model this
+/// functionality should move to clang-sycl-linker or clang-linker-wrapper.
+bool runPreSplitProcessingPipeline(Module &M);
+
 /// Splits the given module \p M according to the given \p Settings.
 Expected<std::vector<SplitModule>>
 splitSYCLModule(std::unique_ptr<Module> M, ModuleSplitterSettings Settings);

llvm/lib/SYCLLowerIR/CleanupSYCLMetadata.cpp

@@ -12,6 +12,11 @@
 //===----------------------------------------------------------------------===//
 
 #include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
+#include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
+#include "llvm/IR/Constants.h"
+#include "llvm/SYCLLowerIR/DeviceGlobals.h"
+#include "llvm/Transforms/Utils/GlobalStatus.h"
 
 #include "llvm/IR/Module.h"
 #include "llvm/Pass.h"
@@ -29,17 +34,113 @@ void cleanupSYCLCompilerMetadata(const Module &M, llvm::StringRef MD) {
   Node->eraseFromParent();
 }
 
+// GV is supposed to be either llvm.compiler.used or llvm.used.
+SmallVector<Constant *>
+eraseGlobalVariableAndReturnOperands(GlobalVariable *GV) {
+  assert(GV->user_empty() && "Users aren't expected");
+  Constant *Initializer = GV->getInitializer();
+  GV->setInitializer(nullptr);
+  GV->eraseFromParent();
+
+  // Destroy the initializer and save operands.
+  SmallVector<Constant *> Operands;
+  Operands.resize(0);
+  for (auto &Op : Initializer->operands())
+    Operands.push_back(cast<Constant>(Op));
+
+  assert(isSafeToDestroyConstant(Initializer) &&
+         "Cannot remove initializer of the given GV");
+
+  Initializer->destroyConstant();
+  return Operands;
+}
+
 } // anonymous namespace
 
 PreservedAnalyses CleanupSYCLMetadataPass::run(Module &M,
                                                ModuleAnalysisManager &MAM) {
   // Remove SYCL module-level metadata that will never be used again to avoid
   // duplication of their operands during llvm-link hence preventing
   // increase of the module size
-  llvm::SmallVector<llvm::StringRef, 2> ModuleMDToRemove = {
-      "sycl_aspects", "sycl_types_that_use_aspects"};
+  SmallVector<StringRef, 2> ModuleMDToRemove = {"sycl_aspects",
+                                                "sycl_types_that_use_aspects"};
   for (const auto &MD : ModuleMDToRemove)
     cleanupSYCLCompilerMetadata(M, MD);
 
   return PreservedAnalyses::all();
 }
+
+PreservedAnalyses
+CleanupSYCLMetadataFromLLVMUsed::run(Module &M, ModuleAnalysisManager &) {
+  GlobalVariable *GV = M.getGlobalVariable("llvm.used");
+  if (!GV)
+    return PreservedAnalyses::all();
+
+  SmallVector<Constant *, 8> IOperands =
+      eraseGlobalVariableAndReturnOperands(GV);
+  // Erase all operands.
+  for (auto *Op : IOperands) {
+    auto StrippedOp = Op->stripPointerCasts();
+    auto *F = dyn_cast<Function>(StrippedOp);
+    if (isSafeToDestroyConstant(Op))
+      (Op)->destroyConstant();
+    else if (F && F->getCallingConv() == CallingConv::SPIR_KERNEL &&
+             !F->use_empty()) {
+      // The element in "llvm.used" array has other users. That is Ok for
+      // specialization constants, but is wrong for kernels.
+      report_fatal_error("Unexpected usage of SYCL kernel");
+    }
+
+    // Remove unused kernel declarations to avoid LLVM IR check fails.
+    if (F && F->isDeclaration() && F->use_empty())
+      F->eraseFromParent();
+  }
+
+  return PreservedAnalyses::none();
+}
+
+PreservedAnalyses
+RemoveDeviceGlobalFromLLVMCompilerUsed::run(Module &M,
+                                            ModuleAnalysisManager &) {
+  GlobalVariable *GV = M.getGlobalVariable("llvm.compiler.used");
+  if (!GV)
+    return PreservedAnalyses::all();
+
+  const auto *VAT = cast<ArrayType>(GV->getValueType());
+  // Destroy the initializer. Keep the operands so we keep the ones we need.
+  SmallVector<Constant *> IOperands = eraseGlobalVariableAndReturnOperands(GV);
+
+  // Iterate through all operands. If they are device_global then we drop them
+  // and erase them if they have no uses afterwards. All other values are kept.
+  SmallVector<Constant *> NewOperands;
+  for (auto *Op : IOperands) {
+    auto *DG = dyn_cast<GlobalVariable>(Op->stripPointerCasts());
+
+    // If it is not a device_global we keep it.
+    if (!DG || !isDeviceGlobalVariable(*DG)) {
+      NewOperands.push_back(Op);
+      continue;
+    }
+
+    // Destroy the device_global operand.
+    if (isSafeToDestroyConstant(Op))
+      Op->destroyConstant();
+
+    // Remove device_global if it no longer has any uses.
+    if (!DG->isConstantUsed())
+      DG->eraseFromParent();
+  }
+
+  // If we have any operands left from the original llvm.compiler.used we create
+  // a new one with the new size.
+  if (!NewOperands.empty()) {
+    ArrayType *ATy = ArrayType::get(VAT->getElementType(), NewOperands.size());
+    GlobalVariable *NGV =
+        new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage,
+                           ConstantArray::get(ATy, NewOperands), "");
+    NGV->setName("llvm.compiler.used");
+    NGV->setSection("llvm.metadata");
+  }
+
+  return PreservedAnalyses::none();
+}

llvm/lib/SYCLLowerIR/ESIMD/ESIMDUtils.cpp

@@ -82,6 +82,12 @@ bool isKernel(const Function &F) {
 
 bool isESIMDKernel(const Function &F) { return isKernel(F) && isESIMD(F); }
 
+bool moduleContainsInvokeSimdBuiltin(Module &M) {
+  return std::any_of(M.begin(), M.end(), [](Function &F) {
+    return !F.isDeclaration() && F.getName().starts_with(INVOKE_SIMD_PREF);
+  });
+}
+
 Type *getVectorTyOrNull(StructType *STy) {
   Type *Res = nullptr;
   while (STy && (STy->getStructNumElements() == 1)) {

llvm/lib/SYCLLowerIR/ModuleSplitter.cpp

@@ -23,10 +23,15 @@
 #include "llvm/IR/PassManager.h"
 #include "llvm/IR/PassManagerImpl.h"
 #include "llvm/IRPrinter/IRPrintingPasses.h"
+#include "llvm/SYCLLowerIR/CleanupSYCLMetadata.h"
+#include "llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h"
 #include "llvm/SYCLLowerIR/DeviceGlobals.h"
+#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerInvokeSimd.h"
 #include "llvm/SYCLLowerIR/SYCLDeviceLibReqMask.h"
+#include "llvm/SYCLLowerIR/SYCLJointMatrixTransform.h"
 #include "llvm/SYCLLowerIR/SYCLUtils.h"
+#include "llvm/SYCLLowerIR/SanitizerKernelMetadata.h"
 #include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Error.h"
@@ -1420,6 +1425,47 @@ Expected<std::vector<SplitModule>> parseSplitModulesFromFile(StringRef File) {
   return Modules;
 }
 
+bool runPreSplitProcessingPipeline(Module &M) {
+  ModulePassManager MPM;
+  ModuleAnalysisManager MAM;
+  MAM.registerPass([&] { return PassInstrumentationAnalysis(); });
+
+  // Propagate ESIMD attribute to wrapper functions to prevent
+  // spurious splits and kernel link errors.
+  MPM.addPass(SYCLFixupESIMDKernelWrapperMDPass());
+
+  // After linking device bitcode "llvm.used" holds references to the kernels
+  // that are defined in the device image. But after splitting device image into
+  // separate kernels we may end up with having references to kernel declaration
+  // originating from "llvm.used" in the IR that is passed to llvm-spirv tool,
+  // and these declarations cause an assertion in llvm-spirv. To workaround this
+  // issue remove "llvm.used" from the input module before performing any other
+  // actions.
+  MPM.addPass(CleanupSYCLMetadataFromLLVMUsed());
+
+  // There may be device_global variables kept alive in "llvm.compiler.used"
+  // to keep the optimizer from wrongfully removing them. llvm.compiler.used
+  // symbols are usually removed at backend lowering, but this is handled here
+  // for SPIR-V since SYCL compilation uses llvm-spirv, not the SPIR-V backend.
+  if (M.getTargetTriple().find("spir") != std::string::npos)
+    MPM.addPass(RemoveDeviceGlobalFromLLVMCompilerUsed());
+
+  // Sanitizer specific passes.
+  if (sycl::isModuleUsingAsan(M) || sycl::isModuleUsingMsan(M) ||
+      sycl::isModuleUsingTsan(M))
+    MPM.addPass(SanitizerKernelMetadataPass());
+
+  // Transform Joint Matrix builtin calls to align them with SPIR-V friendly
+  // LLVM IR specification.
+  MPM.addPass(SYCLJointMatrixTransformPass());
+
+  // Do invoke_simd processing before splitting because this:
+  // - saves processing time (the pass is run once, even though on larger IR)
+  // - doing it before SYCL/ESIMD splitting is required for correctness
+  MPM.addPass(SYCLLowerInvokeSimdPass());
+  return !MPM.run(M, MAM).areAllPreserved();
+}
+
 Expected<std::vector<SplitModule>>
 splitSYCLModule(std::unique_ptr<Module> M, ModuleSplitterSettings Settings) {
   ModuleDesc MD = std::move(M); // makeModuleDesc() ?