[SYCL][ESIMD] Implement IR pass to lower C++ ESIMD intrinsics.

llvm/include/llvm/Demangle/ItaniumDemangle.h

@@ -445,6 +445,14 @@ class EnableIfAttr : public Node {
   }
 };
 
+#ifdef _MSC_VER
+// Workaround for MSVC++ bug (Version 2017, 15.8.9) - w/o this forward
+// declaration, the friend declaration in ObjCProtoName below has no effect
+// and leads to compilation error when ObjCProtoName::Protocol private field
+// is accessed in PointerType::printLeft.
+class PointerType;
+#endif // _MSC_VER
+
 class ObjCProtoName : public Node {
   const Node *Ty;
   StringView Protocol;

llvm/include/llvm/InitializePasses.h

@@ -414,6 +414,7 @@ void initializeStripNonLineTableDebugInfoPass(PassRegistry&);
 void initializeStripSymbolsPass(PassRegistry&);
 void initializeStructurizeCFGPass(PassRegistry&);
 void initializeSYCLLowerWGScopeLegacyPassPass(PassRegistry &);
+void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);
 void initializeTailCallElimPass(PassRegistry&);
 void initializeTailDuplicatePass(PassRegistry&);
 void initializeTargetLibraryInfoWrapperPassPass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

@@ -37,6 +37,7 @@
 #include "llvm/CodeGen/Passes.h"
 #include "llvm/IR/Function.h"
 #include "llvm/IR/IRPrintingPasses.h"
+#include "llvm/SYCLLowerIR/LowerESIMD.h"
 #include "llvm/SYCLLowerIR/LowerWGScope.h"
 #include "llvm/Support/Valgrind.h"
 #include "llvm/Transforms/AggressiveInstCombine/AggressiveInstCombine.h"
@@ -201,6 +202,7 @@ namespace {
       (void) llvm::createMergeICmpsLegacyPass();
       (void) llvm::createExpandMemCmpPass();
       (void)llvm::createSYCLLowerWGScopePass();
+      (void)llvm::createSYCLLowerESIMDPass();
       std::string buf;
       llvm::raw_string_ostream os(buf);
       (void) llvm::createPrintModulePass(os);

llvm/include/llvm/SYCLLowerIR/LowerESIMD.h

@@ -0,0 +1,39 @@
+//===---- LowerESIMD.h - lower Explicit SIMD (ESIMD) constructs -----------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Lowers CM-specific LLVM IR constructs coming out of the front-end. These are:
+// - ESIMD intrinsics, e.g.:
+//     template <typename Ty, int N, int M, int VStride, int Width,
+//       int Stride, int ParentWidth = 0>
+//       sycl::intel::gpu::vector_type_t<Ty, M>
+//       __esimd_rdregion(sycl::intel::gpu::vector_type_t<Ty, N> Input,
+//         uint16_t Offset);
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_SYCLLOWERIR_LOWERESIMD_H
+#define LLVM_SYCLLOWERIR_LOWERESIMD_H
+
+#include "llvm/IR/Function.h"
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+/// SPIRV (ESIMD) target specific pass to transform ESIMD specific constructs
+/// like intrinsics to a form parsable by the ESIMD-aware SPIRV translator.
+class SYCLLowerESIMDPass : public PassInfoMixin<SYCLLowerESIMDPass> {
+public:
+  PreservedAnalyses run(Function &F, FunctionAnalysisManager &,
+                        SmallPtrSet<Type *, 4> &GVTS);
+};
+
+FunctionPass *createSYCLLowerESIMDPass();
+void initializeSYCLLowerESIMDLegacyPassPass(PassRegistry &);
+
+} // namespace llvm
+
+#endif // LLVM_SYCLLOWERIR_LOWERESIMD_H

llvm/lib/SYCLLowerIR/CMakeLists.txt

@@ -1,9 +1,54 @@
+# Lowering of SYCL ESIMD kernels depends on vc-intrinsics
+# NOTE: could have been added earlier from llvm/projects
+if (NOT TARGET LLVMGenXIntrinsics)
+  if (NOT DEFINED LLVMGenXIntrinsics_SOURCE_DIR)
+    message(STATUS "vc-intrinsics are missing. Will try to download them from github.com")
+
+    include(FetchContent)
+    FetchContent_Declare(vc-intrinsics
+      GIT_REPOSITORY https://github.com/intel/vc-intrinsics.git
+      GIT_TAG        cce6e48c28eb850d7dadd30841c0d95f009bbca1
+    )
+    FetchContent_MakeAvailable(vc-intrinsics)
+    FetchContent_GetProperties(vc-intrinsics)
+
+    set(LLVMGenXIntrinsics_SOURCE_DIR ${vc-intrinsics_SOURCE_DIR})
+    set(LLVMGenXIntrinsics_BINARY_DIR ${vc-intrinsics_BINARY_DIR})
+  else()
+    # -DLLVMGenXIntrinsics_SOURCE_DIR is provided
+    message(STATUS "vc-intrinsics are added manually ${LLVMGenXIntrinsics_SOURCE_DIR}")
+
+    set(LLVMGenXIntrinsics_BINARY_DIR ${CMAKE_BINARY_DIR}/vc-intrinsics-build)
+    add_subdirectory(${LLVMGenXIntrinsics_SOURCE_DIR} ${LLVMGenXIntrinsics_BINARY_DIR})
+  endif()
+
+  target_include_directories(LLVMGenXIntrinsics
+    PUBLIC $<BUILD_INTERFACE:${LLVMGenXIntrinsics_SOURCE_DIR}/GenXIntrinsics/include>
+    PUBLIC $<BUILD_INTERFACE:${LLVMGenXIntrinsics_BINARY_DIR}/GenXIntrinsics/include>
+  )
+endif()
+
 add_llvm_component_library(LLVMSYCLLowerIR
   LowerWGScope.cpp
+  LowerESIMD.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR
+  ${LLVM_MAIN_SRC_DIR}/projects/vc-intrinsics/GenXIntrinsics/include
+  ${LLVM_BINARY_DIR}/projects/vc-intrinsics/GenXIntrinsics/include
 
   DEPENDS
   intrinsics_gen
+  LLVMGenXIntrinsics
+  LLVMDemangle
+  LLVMTransformUtils
+
+  LINK_LIBS
+  LLVMGenXIntrinsics
+  LLVMDemangle
+  LLVMTransformUtils
   )
+
+target_include_directories(LLVMSYCLLowerIR
+  PRIVATE ${LLVM_MAIN_SRC_DIR}/projects/vc-intrinsics/GenXIntrinsics/include
+  PRIVATE ${LLVM_BINARY_DIR}/projects/vc-intrinsics/GenXIntrinsics/include)