[ESIMD] Move ESIMD-specific passes to sycl-post-link (#3222)

This patch is a continuation of prior series of patches to allow
having ESIMD and regular SYCL kernels in the same source code
and in the same program.

To mix two types of kernels we need to postpone lowering of
ESIMD specific constructs and do that after we split ESIMD from
regular SYCL code, which is done in sycl-post-link.

This is a finalizing patch that only flips the switch for
doing ESIMD-specific transformations in sycl-post-link
and removes those passes from the common optimization
pipeline in FE (BackendUtil.cpp). All the preliminary patches
have already been committed.

Author: DenisBakhvalov

clang/lib/CodeGen/BackendUtil.cpp

@@ -28,7 +28,6 @@
 #include "llvm/CodeGen/RegAllocRegistry.h"
 #include "llvm/CodeGen/SchedulerRegistry.h"
 #include "llvm/CodeGen/TargetSubtargetInfo.h"
-#include "llvm/GenXIntrinsics/GenXSPIRVWriterAdaptor.h"
 #include "llvm/IR/DataLayout.h"
 #include "llvm/IR/IRPrintingPasses.h"
 #include "llvm/IR/LegacyPassManager.h"
@@ -42,7 +41,6 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
-#include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/MemoryBuffer.h"
@@ -839,25 +837,6 @@ void EmitAssemblyHelper::CreatePasses(legacy::PassManager &MPM,
 
   PMBuilder.populateFunctionPassManager(FPM);
   PMBuilder.populateModulePassManager(MPM);
-
-  // Customize the tail of the module passes list for the ESIMD extension.
-  if (LangOpts.SYCLIsDevice && LangOpts.SYCLExplicitSIMD &&
-      CodeGenOpts.OptimizationLevel != 0) {
-    MPM.add(createESIMDLowerVecArgPass());
-    MPM.add(createESIMDLowerLoadStorePass());
-    MPM.add(createSROAPass());
-    MPM.add(createEarlyCSEPass(true));
-    MPM.add(createInstructionCombiningPass());
-    MPM.add(createDeadCodeEliminationPass());
-    MPM.add(createFunctionInliningPass(
-        CodeGenOpts.OptimizationLevel, CodeGenOpts.OptimizeSize,
-        (!CodeGenOpts.SampleProfileFile.empty() &&
-         CodeGenOpts.PrepareForThinLTO)));
-    MPM.add(createSROAPass());
-    MPM.add(createEarlyCSEPass(true));
-    MPM.add(createInstructionCombiningPass());
-    MPM.add(createDeadCodeEliminationPass());
-  }
 }
 
 static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) {
@@ -954,11 +933,6 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
   PerFunctionPasses.add(
       createTargetTransformInfoWrapperPass(getTargetIRAnalysis()));
 
-  // ESIMD extension always requires lowering of certain IR constructs, such as
-  // ESIMD C++ intrinsics, as the last FE step.
-  if (LangOpts.SYCLIsDevice && LangOpts.SYCLExplicitSIMD)
-    PerModulePasses.add(createSYCLLowerESIMDPass());
-
   CreatePasses(PerModulePasses, PerFunctionPasses);
 
   legacy::PassManager CodeGenPasses;
@@ -976,9 +950,6 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action,
       !LangOpts.SYCLExplicitSIMD && LangOpts.EnableDAEInSpirKernels)
     PerModulePasses.add(createDeadArgEliminationSYCLPass());
 
-  if (LangOpts.SYCLIsDevice && LangOpts.SYCLExplicitSIMD)
-    PerModulePasses.add(createGenXSPIRVWriterAdaptorPass());
-
   switch (Action) {
   case Backend_EmitNothing:
     break;

clang/lib/CodeGen/CMakeLists.txt

@@ -27,14 +27,6 @@ set(LLVM_LINK_COMPONENTS
   TransformUtils
   )
 
-
-get_property(LLVMGenXIntrinsics_SOURCE_DIR GLOBAL PROPERTY LLVMGenXIntrinsics_SOURCE_PROP)
-get_property(LLVMGenXIntrinsics_BINARY_DIR GLOBAL PROPERTY LLVMGenXIntrinsics_BINARY_PROP)
-
-include_directories(
-  ${LLVMGenXIntrinsics_SOURCE_DIR}/GenXIntrinsics/include
-  ${LLVMGenXIntrinsics_BINARY_DIR}/GenXIntrinsics/include)
-
 add_clang_library(clangCodeGen
   BackendUtil.cpp
   CGAtomic.cpp
@@ -96,14 +88,9 @@ add_clang_library(clangCodeGen
   TargetInfo.cpp
   VarBypassDetector.cpp
 
-  ADDITIONAL_HEADER_DIRS
-  ${LLVMGenXIntrinsics_SOURCE_DIR}/GenXIntrinsics/include
-  ${LLVMGenXIntrinsics_BINARY_DIR}/GenXIntrinsics/include
-
   DEPENDS
   ${codegen_deps}
   intrinsics_gen
-  LLVMGenXIntrinsics
 
   LINK_LIBS
   clangAnalysis
@@ -112,5 +99,4 @@ add_clang_library(clangCodeGen
   clangFrontend
   clangLex
   clangSerialization
-  LLVMGenXIntrinsics
   )

clang/lib/Driver/ToolChains/Clang.cpp

@@ -8348,7 +8348,7 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA,
                        options::OPT_fno_sycl_device_code_split_esimd, true))
       addArgs(CmdArgs, TCArgs, {"-split-esimd"});
     if (TCArgs.hasFlag(options::OPT_fsycl_device_code_lower_esimd,
-                       options::OPT_fno_sycl_device_code_lower_esimd, false))
+                       options::OPT_fno_sycl_device_code_lower_esimd, true))
       addArgs(CmdArgs, TCArgs, {"-lower-esimd"});
   }
   addArgs(CmdArgs, TCArgs,

clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp

@@ -1,6 +1,13 @@
+// TODO: previously code generation and ESIMD lowering was
+// a part of the same %clang_cc1 invocation, but now it is
+// separate. So, we can split this test into 2, where one
+// will be testing code generation and the second ESIMD lowering.
+//
 // RUN: %clang_cc1 -fsycl -fsycl-explicit-simd -fsycl-is-device \
 // RUN:   -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice \
-// RUN:   -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+// RUN:   -disable-llvm-passes -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // This test checks
 // 1) proper metadata generation for accessors used in ESIMD

clang/test/CodeGenSYCL/esimd-private-global.cpp

@@ -1,14 +1,27 @@
+// TODO: previously code generation and ESIMD lowering was
+// a part of the same %clang_cc1 invocation, but now it is
+// separate. So, we can split this test into 2, where one
+// will be testing code generation and the second ESIMD lowering.
+//
 // RUN: %clang_cc1 -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice \
-// RUN:   -fsycl -fsycl-is-device -fsycl-explicit-simd -emit-llvm %s -o - | \
-// RUN:   FileCheck %s
+// RUN:   -fsycl -fsycl-is-device -fsycl-explicit-simd -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // This test checks that FE allows globals with register_num attribute in ESIMD mode.
 
 __attribute__((opencl_private)) __attribute__((register_num(17))) int vc;
 // CHECK: @vc = {{.+}} i32 0, align 4 #0
 
-SYCL_EXTERNAL void init_vc(int x) {
-  vc = x;
-  // CHECK: store i32 %{{[0-9a-zA-Z_]+}}, i32* @vc
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+void init_vc(int x) {
+  kernel<class kernel_esimd>([=]() __attribute__((sycl_explicit_simd)) {
+    vc = x;
+    // CHECK: store i32 %{{[0-9a-zA-Z_]+}}, i32* @vc
+  });
 }
 // CHECK: attributes #0 = {{.*"VCByteOffset"="17".*"VCVolatile"}}

clang/test/Driver/sycl-device-lib.cpp

@@ -121,7 +121,7 @@
 // RUN:   | FileCheck %s -check-prefix=SYCL_LLVM_LINK_NO_DEVICE_LIB
 // SYCL_LLVM_LINK_NO_DEVICE_LIB: clang{{.*}} "-cc1" {{.*}} "-fsycl-is-device"
 // SYCL_LLVM_LINK_NO_DEVICE_LIB-NOT: llvm-link{{.*}}  "-only-needed"
-// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-split-esimd" "-O2" "-spec-const=rt" "-o" "{{.*}}.table" "{{.*}}.bc"
+// SYCL_LLVM_LINK_NO_DEVICE_LIB: sycl-post-link{{.*}}  "-symbols" "-split-esimd" "-lower-esimd" "-O2" "-spec-const=rt" "-o" "{{.*}}.table" "{{.*}}.bc"
 
 /// ###########################################################################
 /// test llvm-link behavior for special user input whose filename resembles SYCL device library

clang/test/Driver/sycl-offload-with-split.c

@@ -318,8 +318,8 @@
 // CHK-NO-ESIMD-SPLIT-NOT: sycl-post-link{{.*}} "-split-esimd"
 
 // Check lowering of ESIMD device code.
-// RUN:   %clang    -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-LOWER
-// RUN:   %clang_cl -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-LOWER
+// RUN:   %clang    -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
+// RUN:   %clang_cl -### -fsycl %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang    -### -fsycl -fsycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang_cl -### -fsycl -fsycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-ESIMD-LOWER
 // RUN:   %clang    -### -fsycl -fno-sycl-device-code-lower-esimd %s 2>&1 | FileCheck %s -check-prefixes=CHK-NO-ESIMD-LOWER

sycl/test/esimd/glob.cpp

@@ -1,5 +1,6 @@
-// RUN: %clangxx -fsycl -fsycl-explicit-simd -c -fsycl-device-only -Xclang -emit-llvm %s -o - | \
-// RUN:  FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-explicit-simd -c -fsycl-device-only -Xclang -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // This test checks that globals with register attribute are allowed in ESIMD
 // mode, can be accessed in functions and correct LLVM IR is generated
@@ -9,7 +10,6 @@
 #include <CL/sycl/INTEL/esimd.hpp>
 #include <iostream>
 
-using namespace cl::sycl;
 using namespace sycl::INTEL::gpu;
 
 constexpr unsigned VL = 16;
@@ -22,7 +22,16 @@ ESIMD_PRIVATE ESIMD_REGISTER(17 + VL) simd<int, VL> vc1;
 // CHECK-DAG: @vc1 = {{.+}} <16 x i32> zeroinitializer, align 64 #1
 // CHECK-DAG: attributes #1 = { {{.*}}"VCByteOffset"="33" "VCGlobalVariable" "VCVolatile"{{.*}} }
 
-SYCL_EXTERNAL ESIMD_NOINLINE void init_vc(int x) {
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+SYCL_ESIMD_FUNCTION SYCL_EXTERNAL ESIMD_NOINLINE void init_vc(int x) {
   vc1 = vc + 1;
   vc = x;
 }
+
+void caller(int x) {
+  kernel<class kernel_esimd>([=]() SYCL_ESIMD_KERNEL { init_vc(x); });
+}

sycl/test/esimd/intrins_trans.cpp

@@ -1,5 +1,6 @@
-// RUN: %clangxx -O0 -fsycl -fsycl-explicit-simd -fsycl-device-only -Xclang -emit-llvm %s -o - | \
-// RUN:   FileCheck %s
+// RUN: %clangxx -O0 -fsycl -fsycl-explicit-simd -fsycl-device-only -Xclang -emit-llvm %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
 
 // Checks ESIMD intrinsic translation.
 // NOTE: must be run in -O0, as optimizer optimizes away some of the code
@@ -109,7 +110,7 @@ SYCL_ESIMD_FUNCTION SYCL_EXTERNAL simd<float, 16> foo() {
   // CHECK: call void @llvm.genx.media.st.v32i32(i32 0, i32 %[[SI2]], i32 0, i32 32, i32 %{{[0-9a-zA-Z_.]+}}, i32 %{{[0-9a-zA-Z_.]+}}, <32 x i32> %{{[0-9a-zA-Z_.]+}})
 
   auto ee = __esimd_vload<int, 16>((vector_type_t<int, 16> *)(&vg));
-  // CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.vload.v16i32.p4v16i32(<16 x i32> addrspace(4)* {{.*}})
+  // CHECK: %{{[0-9a-zA-Z_.]+}} = call <16 x i32> @llvm.genx.vload.v16i32.p0v16i32(<16 x i32>* {{.*}})
   __esimd_vstore<int, 32>(&vc, va.data());
   // CHECK: store <32 x i32>  %{{[0-9a-zA-Z_.]+}}, <32 x i32> addrspace(4)* {{.*}}
 

sycl/test/esimd/spirv_intrins_trans.cpp

@@ -1,6 +1,9 @@
-// RUN: %clangxx -fsycl -fsycl-explicit-simd -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-explicit-simd -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o %t
+// RUN: sycl-post-link -split-esimd -lower-esimd -O2 -S %t -o %t.table
+// RUN: FileCheck %s -input-file=%t_esimd_0.ll
+
 // This test checks that all SPIR-V intrinsics are correctly
-// translated into GenX counterparts (implemented in LowerCM.cpp)
+// translated into GenX counterparts (implemented in LowerESIMD.cpp)
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/INTEL/esimd.hpp>