[SYCL][ESIMD] Add hidden compiler option to not compile host code (#11037)

ESIMD code can only be run on an Intel GPU. Previously we had the ESMID
emulator that called into host implementations of some of the GPU
intrinsics, but that has been deprecated and will be removed. Even when
the emulator was supported, it was seldom used by customers.

To summarize, the host compile for ESIMD code had little value to
customers previous, and will soon have zero value once the emulator is
removed.

However, the host compile can be slow. We have a bug report from a
customer reporting slow compilation time for an ESIMD kernel with a
large vector. The reason for the slowdown was investigated by the X86
backend team and they found for the host compile, the vector exceeds the
maximum legal vector size on X86 and instruction selection needs to
spend a lot of time doing heavy lifting to generate the right assembly.

This is a total waste of the user's time as that code will never and can
never be run.

As a workaround for these cases, add a new hidden option,
`-fsycl-esimd-build-host-code`/`-fno-sycl-esimd-build-host-code`,
default `-fsycl-esimd-build-host-code`, to try to speed up the host
compile. This is mostly done through a new pass which runs early and
removes all IR for ESIMD functions and replaces them with a single
return statement. This allows later optimizations to do even more
cleanup.

Also, I modified some of the ESIMD macros to generate IR that is more
likely to be optimized out for the host, just disabling inlining and
setting internal linkage.

We do not intend to advertise this option as there should be no issue in
most cases, but we would like to be able to provide a workaround if
users hit an issue.

I have manually verified this option fixes the slowdown reported by the
customer.

---------

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

Author: sarnex

clang/include/clang/Basic/LangOptions.def

@@ -281,6 +281,7 @@ LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for
 LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
 LANGOPT(SYCLForceInlineKernelLambda , 1, 0, "Force inline SYCL kernel lambdas in entry point")
 LANGOPT(SYCLESIMDForceStatelessMem, 1, 0, "Make accessors use USM memory in ESIMD kernels")
+LANGOPT(SYCLESIMDBuildHostCode, 1, 1, "Build the host implementation of ESIMD functions")
 ENUM_LANGOPT(SYCLVersion  , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used")
 LANGOPT(DeclareSPIRVBuiltins, 1, 0, "Declare SPIR-V builtin functions")
 LANGOPT(SYCLExplicitSIMD  , 1, 0, "SYCL compilation with explicit SIMD extension")

clang/include/clang/Driver/Options.td

@@ -3634,7 +3634,14 @@ defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-me
             "Disabled by default. (experimental)">,
     NegFlag<SetFalse, [], [ClangOption], "Do not enforce using stateless memory accesses. (experimental)">,
     BothFlags<[], [ClangOption, CLOption, DXCOption, CC1Option], "">>;
-
+// TODO: Remove this option once ESIMD headers are updated to
+// guard vectors to be device only.
+defm sycl_esimd_build_host_code : BoolFOption<"sycl-esimd-build-host-code",
+    LangOpts<"SYCLESIMDBuildHostCode">, DefaultTrue,
+    PosFlag<SetTrue, [], [ClangOption], "Build the host implementation of ESIMD functions."
+            "Enabled by default.">,
+    NegFlag<SetFalse, [], [ClangOption], "Don't build the host implementation of ESIMD functions.">,
+    BothFlags<[HelpHidden], [ClangOption, CLOption, DXCOption, CC1Option], "">>;    
 def fsycl_targets_EQ : CommaJoined<["-"], "fsycl-targets=">, Flags<[NoXarchOption]>,
   Visibility<[ClangOption, CLOption, DXCOption, CC1Option]>,
   HelpText<"Specify comma-separated list of triples SYCL offloading targets to be supported">;

clang/lib/CodeGen/BackendUtil.cpp

@@ -45,6 +45,7 @@
 #include "llvm/Passes/StandardInstrumentations.h"
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
+#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
 #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
@@ -965,6 +966,11 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
             MPM.addPass(
                 SYCLPropagateAspectsUsagePass(/*ExcludeAspects=*/{"fp64"}));
           });
+    else if (LangOpts.SYCLIsHost && !LangOpts.SYCLESIMDBuildHostCode)
+      PB.registerPipelineStartEPCallback(
+          [&](ModulePassManager &MPM, OptimizationLevel Level) {
+            MPM.addPass(ESIMDRemoveHostCodePass());
+          });
 
     // Add the InferAddressSpaces pass for all the SPIR[V] targets
     if (TargetTriple.isSPIR() || TargetTriple.isSPIRV()) {

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5429,6 +5429,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     else {
       for (auto &Macro : D.getSYCLTargetMacroArgs())
         CmdArgs.push_back(Args.MakeArgString(Macro));
+      if (!Args.hasFlag(options::OPT_fsycl_esimd_build_host_code,
+                        options::OPT_fno_sycl_esimd_build_host_code, true))
+        CmdArgs.push_back("-fno-sycl-esimd-build-host-code");
     }
     if (Args.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem,
                      options::OPT_fno_sycl_esimd_force_stateless_mem, false))

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1339,6 +1339,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
       Builder.defineMacro("__ENABLE_USM_ADDR_SPACE__");
       Builder.defineMacro("SYCL_DISABLE_FALLBACK_ASSERT");
     }
+  } else if (LangOpts.SYCLIsHost && LangOpts.SYCLESIMDBuildHostCode) {
+    Builder.defineMacro("__ESIMD_BUILD_HOST_CODE");
   }
   if (LangOpts.SYCLUnnamedLambda)
     Builder.defineMacro("__SYCL_UNNAMED_LAMBDA__");

clang/test/Driver/sycl-esimd-build-host-code.cpp

@@ -0,0 +1,13 @@
+
+/// Verify that the driver option is translated to corresponding options
+/// to host
+// RUN: %clang -### -fsycl -fno-sycl-esimd-build-host-code \
+// RUN: %s 2>&1 | FileCheck -check-prefix=CHECK-PASS-TO-COMPS %s
+// CHECK-PASS-TO-COMPS: clang
+// CHECK-PASS-TO-COMPS-NOT: "-fno-sycl-esimd-build-host-code"
+// CHECK-PASS-TO-COMPS: sycl-post-link{{.*}}
+// CHECK-PASS-TO-COMPS: clang{{.*}} "-fsycl-is-host" {{.*}}"-fno-sycl-esimd-build-host-code"
+
+
+/// Verify that removing host code is not enabled by default
+// RUN: %clang -### -fsycl %s 2>&1 | FileCheck -implicit-check-not "-fno-sycl-esimd-build-host-code" %s

llvm/include/llvm/SYCLLowerIR/ESIMD/LowerESIMD.h

@@ -80,6 +80,12 @@ class SYCLFixupESIMDKernelWrapperMDPass
 public:
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
 };
+
+class ESIMDRemoveHostCodePass : public PassInfoMixin<ESIMDRemoveHostCodePass> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &);
+};
+
 } // namespace llvm
 
 #endif // LLVM_SYCLLOWERIR_LOWERESIMD_H

llvm/lib/Passes/PassRegistry.def

@@ -137,6 +137,7 @@ MODULE_PASS("SPIRITTAnnotations", SPIRITTAnnotationsPass())
 MODULE_PASS("deadargelim-sycl", DeadArgumentEliminationSYCLPass())
 MODULE_PASS("sycllowerwglocalmemory", SYCLLowerWGLocalMemoryPass())
 MODULE_PASS("lower-esimd-kernel-attrs", SYCLFixupESIMDKernelWrapperMDPass())
+MODULE_PASS("esimd-remove-host-code", ESIMDRemoveHostCodePass());
 MODULE_PASS("sycl-propagate-aspects-usage", SYCLPropagateAspectsUsagePass())
 MODULE_PASS("sycl-add-opt-level-attribute", SYCLAddOptLevelAttributePass())
 MODULE_PASS("compile-time-properties", CompileTimePropertiesPass())

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -50,6 +50,7 @@ add_llvm_component_library(LLVMSYCLLowerIR
   ESIMD/ESIMDOptimizeVecArgCallConv.cpp
   ESIMD/ESIMDUtils.cpp
   ESIMD/ESIMDVerifier.cpp
+  ESIMD/ESIMDRemoveHostCode.cpp
   ESIMD/LowerESIMD.cpp
   ESIMD/LowerESIMDKernelAttrs.cpp
   CompileTimePropertiesPass.cpp

llvm/lib/SYCLLowerIR/ESIMD/ESIMDRemoveHostCode.cpp

@@ -0,0 +1,75 @@
+//===-- ESIMDRemoveHostCode.cpp - remove host code for ESIMD -----------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// ESIMD code is not often run on the host, but we still compile for the host.
+// If requested by the user, remove the implementations of all ESIMD functions
+// to possibly speed up host compilation time.
+//===----------------------------------------------------------------------===//
+
+#define DEBUG_TYPE "ESIMDRemoveHostCodePass"
+
+#include "llvm/Demangle/Demangle.h"
+#include "llvm/Demangle/ItaniumDemangle.h"
+#include "llvm/IR/Function.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Pass.h"
+#include "llvm/SYCLLowerIR/ESIMD/ESIMDUtils.h"
+#include "llvm/SYCLLowerIR/ESIMD/LowerESIMD.h"
+#include "llvm/Support/Debug.h"
+#include "llvm/TargetParser/Triple.h"
+
+using namespace llvm;
+using namespace llvm::esimd;
+namespace id = itanium_demangle;
+PreservedAnalyses ESIMDRemoveHostCodePass::run(Module &M,
+                                               ModuleAnalysisManager &) {
+  // TODO: Remove this pass once ESIMD headers are updated to
+  // guard vectors to be device only.
+  bool Modified = false;
+  assert(!Triple(M.getTargetTriple()).isSPIR() &&
+         "Pass should not be run for SPIR targets");
+  for (auto &F : M.functions()) {
+    if (F.isDeclaration())
+      continue;
+    StringRef MangledName = F.getName();
+    id::ManglingParser<SimpleAllocator> Parser(MangledName.begin(),
+                                               MangledName.end());
+    id::Node *AST = Parser.parse();
+    if (!AST || AST->getKind() != id::Node::KFunctionEncoding)
+      continue;
+
+    auto *FE = static_cast<id::FunctionEncoding *>(AST);
+    const id::Node *NameNode = FE->getName();
+    if (!NameNode)
+      continue;
+
+    id::OutputBuffer NameBuf;
+    NameNode->print(NameBuf);
+    StringRef Name(NameBuf.getBuffer(), NameBuf.getCurrentPosition());
+    if (!Name.startswith("sycl::_V1::ext::intel::esimd::") &&
+        !Name.startswith("sycl::_V1::ext::intel::experimental::esimd::"))
+      continue;
+    SmallVector<BasicBlock *> BBV;
+    for (BasicBlock &BB : F) {
+      BB.dropAllReferences();
+      BBV.push_back(&BB);
+    }
+    for (auto *BB : BBV)
+      BB->eraseFromParent();
+
+    Value *Ret = nullptr;
+    Type *RetTy = F.getFunctionType()->getReturnType();
+    if (!RetTy->isVoidTy())
+      Ret = Constant::getNullValue(RetTy);
+
+    LLVMContext &Ctx = F.getParent()->getContext();
+    BasicBlock *BB = BasicBlock::Create(Ctx, "", &F);
+    ReturnInst::Create(Ctx, Ret, BB);
+    Modified = true;
+  }
+  return Modified ? PreservedAnalyses::none() : PreservedAnalyses::all();
+}

llvm/test/SYCLLowerIR/ESIMD/remove_host_code.ll

@@ -0,0 +1,32 @@
+; RUN: opt -passes=esimd-remove-host-code -S < %s | FileCheck %s
+
+; This test checks that ESIMDRemoveHostCode removes all code from ESIMD
+; functions and leaves others untouched.
+
+target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
+target triple = "x86_64-unknown-linux-gnu"
+
+; Function Attrs: alwaysinline mustprogress uwtable
+define linkonce_odr dso_local void @foo() {
+; CHECK: foo
+; CHECK-NEXT: %1 = alloca double, align 8
+; CHECK-NEXT: ret void
+%1 = alloca double, align 8
+ret void
+}
+
+; Function Attrs: alwaysinline mustprogress uwtable
+define linkonce_odr dso_local void @_ZN4sycl3_V13ext5intel12experimental5esimd15lsc_block_storeIdLi64ELNS4_13lsc_data_sizeE0ELNS4_10cache_hintE0ELS7_0ENS2_5esimd6detail26dqword_element_aligned_tagEEENSt9enable_ifIXsr4sycl3ext5intel5esimdE19is_simd_flag_type_vIT4_EEvE4typeEPT_NS8_4simdISF_XT0_EEENS9_14simd_mask_implItLi1EEESC_() {
+; CHECK: lsc_block_store
+; CHECK-NEXT: ret void
+%1 = alloca double, align 8
+ret void
+}
+
+; Function Attrs: alwaysinline mustprogress uwtable
+define linkonce_odr dso_local ptr @_ZN4sycl3_V13ext5intel12experimental5esimd15lsc_block_fobarIdLi64ELNS4_13lsc_data_sizeE0ELNS4_10cache_hintE0ELS7_0ENS2_5esimd6detail26dqword_element_aligned_tagEEENSt9enable_ifIXsr4sycl3ext5intel5esimdE19is_simd_flag_type_vIT4_EEvE4typeEPT_NS8_4simdISF_XT0_EEENS9_14simd_mask_implItLi1EEESC_() {
+; CHECK: lsc_block_fobar
+; CHECK-NEXT: ret ptr null
+%1 = alloca double, align 8
+ret ptr %1
+}

sycl/include/sycl/ext/intel/esimd/detail/defines_elementary.hpp

@@ -37,8 +37,11 @@
 // each work-item is mapped to a separate OS thread on host device.
 #define ESIMD_PRIVATE thread_local
 #define ESIMD_REGISTER(n)
-
+#ifdef __ESIMD_BUILD_HOST_CODE
 #define __ESIMD_API ESIMD_INLINE
+#else // __ESIMD_BUILD_HOST_CODE
+#define __ESIMD_API ESIMD_NOINLINE __attribute__((internal_linkage))
+#endif // __ESIMD_BUILD_HOST_CODE
 #endif // __SYCL_DEVICE_ONLY__
 
 // Mark a function being noinline

sycl/include/sycl/ext/intel/esimd/detail/util.hpp

@@ -19,6 +19,8 @@
 
 #ifdef __SYCL_DEVICE_ONLY__
 #define __ESIMD_INTRIN __DPCPP_SYCL_EXTERNAL SYCL_ESIMD_FUNCTION
+#elif !defined(__ESIMD_BUILD_HOST_CODE)
+#define __ESIMD_INTRIN ESIMD_NOINLINE __attribute__((internal_linkage))
 #else
 #define __ESIMD_INTRIN inline
 #endif // __SYCL_DEVICE_ONLY__

sycl/test-e2e/ESIMD/no_host_code.cpp

@@ -0,0 +1,12 @@
+//==---------------- no_host_code.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: esimd_emulator
+// RUN: %{build} -fno-sycl-esimd-build-host-code -o %t.out
+// RUN: %{run} %t.out
+
+#include "BitonicSortK.hpp"