[Conversion][MathToVC] Add MathToVC pass. (#943)

Convert selective math ops to vc-intrinsics based on condition.

Currently supported ops:
math.ceil,
math.floor,
math.exp (only for vectors),
math.exp2 (only for vectors).

The pass provides an option for high-precision interim calculation.
Exp op conversion requires multuplication of input to log2e, this option,
allows these operations to be done in f32 if the input data type is lower precsion.
Default is to do these operation in actual precision of the exp input.

include/imex/Conversion/CMakeLists.txt

@@ -9,3 +9,4 @@ add_subdirectory(DistToStandard)
 add_subdirectory(DropRegions)
 add_subdirectory(XeTileToXeGPU)
 add_subdirectory(XeGPUToVC)
+add_subdirectory(MathToVC)

include/imex/Conversion/MathToVC/MathToVC.h

@@ -0,0 +1,51 @@
+//===- MathToVC.h - Conversion---------------*- C++ -*-===//
+//
+// Copyright 2024 Intel Corporation
+// Part of the IMEX 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file implements conversion of the select math dialect operations into
+/// Func dialect calls to vc-intrinsics functions
+///
+//===----------------------------------------------------------------------===//
+#ifndef IMEX_CONVERSION_MATHTOVC_H
+#define IMEX_CONVERSION_MATHTOVC_H
+
+#include <mlir/Dialect/Math/IR/Math.h>
+#include <mlir/Dialect/Vector/IR/VectorOps.h>
+
+#include "imex/Utils/XeCommon.h"
+
+namespace mlir {
+
+class ConversionTarget;
+class LLVMTypeConverter;
+class Pass;
+class Operation;
+class RewritePatternSet;
+template <typename T> class OperationPass;
+
+namespace gpu {
+class GPUModuleOp;
+} // namespace gpu
+
+} // namespace mlir
+
+namespace imex {
+#define GEN_PASS_DECL_CONVERTMATHTOVC
+#include "imex/Conversion/Passes.h.inc"
+
+void populateMathToVCPatterns(
+    ::mlir::LLVMTypeConverter &typeConverter,
+    ::mlir::RewritePatternSet &patterns,
+    bool enableHighPrecisionInterimCalculation = false);
+void configureMathToVCConversionLegality(::mlir::ConversionTarget &target);
+std::unique_ptr<::mlir::OperationPass<::mlir::gpu::GPUModuleOp>>
+createConvertMathToVCPass();
+
+} // namespace imex
+#endif

include/imex/Conversion/Passes.h

@@ -22,6 +22,7 @@
 #include <imex/Conversion/GPUToGPUX/GPUToGPUX.h>
 #include <imex/Conversion/GPUToSPIRV/GPUToSPIRVPass.h>
 #include <imex/Conversion/GPUXToLLVM/GPUXToLLVMPass.h>
+#include <imex/Conversion/MathToVC/MathToVC.h>
 #include <imex/Conversion/NDArrayToLinalg/NDArrayToLinalg.h>
 #include <imex/Conversion/XeGPUToVC/XeGPUToVC.h>
 #include <imex/Conversion/XeTileToXeGPU/XeTileToXeGPU.h>

include/imex/Conversion/Passes.td

@@ -427,4 +427,34 @@ def ConvertXeGPUToVC : Pass<"convert-xegpu-to-vc", "::mlir::gpu::GPUModuleOp"> {
   let constructor = "imex::createConvertXeGPUToVCPass()";
 }
 
+
+//===----------------------------------------------------------------------===//
+// MathToVC
+//===----------------------------------------------------------------------===//
+// high-precision-interim-calculation
+def ConvertMathToVC : Pass<"convert-math-to-vc", "::mlir::gpu::GPUModuleOp"> {
+  let summary = "Generate vc-intrinsics functions for select math dialect operations";
+  let description = [{
+    Convert select math dialect operations into the Func dialect calls to vc-intrinsics
+    functions.
+    Some math operations are not supported by the VC compiler (IGC vector backend)
+    and need to be converted to vc-intrinsic calls.
+    This pass converts these math operations to vc-intrinsics.
+    }];
+  let options = [
+    Option<"enableHighPrecisionInterimCalculation", "enable-high-precision-interim-calculation", "bool",
+           /*default=*/"false",
+           "Enables high precision (f32) interim  calculation for math operations."
+           "For any interim instruction added as part of the conversion will be high precision(f32).">
+  ];
+
+  let dependentDialects = ["::mlir::math::MathDialect",
+                           "::mlir::vector::VectorDialect",
+                            "::mlir::LLVM::LLVMDialect",
+                            "::mlir::func::FuncDialect",
+                           "::mlir::arith::ArithDialect"
+                            ];
+  let constructor = "imex::createConvertMathToVCPass()";
+}
+
 #endif // _IMEX_CONVERSION_PASSES_TD_INCLUDED_

lib/Conversion/XeGPUToVC/Utils.h → include/imex/Utils/VCUtils.h

@@ -8,14 +8,20 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// This file defines some utils used in XeGPUToVC pass
+/// This file defines some utils used in ConversionToVC passes
 ///
 //===----------------------------------------------------------------------===//
 
-#ifndef XEGPU_VC_UTILS_H
-#define XEGPU_VC_UTILS_H
+#ifndef VC_UTILS_H
+#define VC_UTILS_H
 
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
+#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
 #include "mlir/IR/Builders.h"
+#include "mlir/Transforms/DialectConversion.h"
 
 using namespace mlir;
 

include/imex/Utils/XeCommon.h

@@ -617,6 +617,12 @@ llvm::SmallVector<T> swapLastTwoElements(llvm::ArrayRef<T> shape) {
 ///   output strides = 60x20x5x1
 llvm::SmallVector<int64_t> defaultStrides(llvm::ArrayRef<int64_t> shape);
 
+/// Checks if the given `type` is a 1-D vector type that requires VectorAnyINTEL
+/// capability. In other words, the vector size is not supported by SPIR-V.
+/// SPIR-V only supports 2, 3, 4, 8, 16 elements (8 and 16 with Vector16
+/// capability).
+bool isVectorAnyINTELType(mlir::Type type);
+
 } // namespace imex
 
 #endif

lib/Conversion/CMakeLists.txt

@@ -4,5 +4,6 @@ add_subdirectory(DropRegions)
 add_subdirectory(GPUToSPIRV)
 add_subdirectory(GPUToGPUX)
 add_subdirectory(GPUXToLLVM)
+add_subdirectory(MathToVC)
 add_subdirectory(XeTileToXeGPU)
 add_subdirectory(XeGPUToVC)

lib/Conversion/MathToVC/CMakeLists.txt

@@ -0,0 +1,21 @@
+add_imex_conversion_library(IMEXMathToVC
+  MathToVC.cpp
+
+
+  ADDITIONAL_HEADER_DIRS
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Conversion/MathToVC
+
+  DEPENDS
+  IMEXConversionPassIncGen
+
+  #LINK_COMPONENTS
+
+  LINK_LIBS PUBLIC
+  MLIRIR
+  MLIRSupport
+  # MLIRTransforms
+  MLIRLLVMCommonConversion
+
+  MLIRGPUDialect
+  MLIRPass
+  )