Add XeGPUToVC pass to lower XeGPU ops to VC intrinsics Func calls (intel#720)

This change adds following:
            - Adds a new pass xegpu-to-vc to lower XeGPU dialect ops
              to vc-intrinsics functions via func dialect. IR generated from this

            - [xegpu][func][spirv] Add necessary linkage attributes to intrinsics needed for SPIR-V dialect path
            - [xegpu][func][spirv] Add VectorComputeFunctionINTEL attribute to intrinsic functions

Co-authored-by: Md Abdullah Shahneous Bari <[email protected]>

Author: drprajap

include/imex/Conversion/CMakeLists.txt

@@ -8,3 +8,4 @@ add_mlir_doc(Passes IMEXConversionPasses ./ -gen-pass-doc)
 add_subdirectory(DistToStandard)
 add_subdirectory(DropRegions)
 add_subdirectory(XeTileToXeGPU)
+add_subdirectory(XeGPUToVC)

include/imex/Conversion/Passes.h

@@ -24,6 +24,7 @@
 #include <imex/Conversion/GPUXToLLVM/GPUXToLLVMPass.h>
 #include <imex/Conversion/NDArrayToLinalg/NDArrayToLinalg.h>
 #include <imex/Conversion/XeGPUToSPIRV/XeGPUToSPIRV.h>
+#include <imex/Conversion/XeGPUToVC/XeGPUToVC.h>
 #include <imex/Conversion/XeTileToXeGPU/XeTileToXeGPU.h>
 
 namespace imex {

include/imex/Conversion/Passes.td

@@ -410,4 +410,36 @@ def ConvertXeTileToXeGPU: Pass<"convert-xetile-to-xegpu", "::mlir::gpu::GPUModul
  ];
 }
 
+
+//===----------------------------------------------------------------------===//
+// XeGPUToVC
+//===----------------------------------------------------------------------===//
+
+def ConvertXeGPUToVC : Pass<"convert-xegpu-to-vc", "::mlir::gpu::GPUModuleOp"> {
+  let summary = "Generate vc-intrinsics functions for xegpu dialect operations";
+  let description = [{
+    Convert XeGPU dialect operations into the Func dialect calls to vc-intrinsics
+    functions.
+    }];
+  let options = [
+    Option<"enableJointMatrix", "enable-joint-matrix","bool", "false",
+          "Enable XeGPU SIMT mode Ops lowered to JointMatrix based Ops">,
+    Option<"enableGenISAIntrinsic", "enable-genisa-intrinsic","bool", "false",
+           "Enable XeGPU SIMT mode Ops lowered to JointMatrix based Ops">,
+    Option<"enableVCIntrinsic", "enable-vc-intrinsic","bool", "true",
+           "Enable XeGPU Ops lowered to intel vc Intrinsics">,
+    Option<"useRawSend", "useRawSend","bool", "true",
+           "Prefer Raw Send API for load/store">
+  ];
+
+  let dependentDialects = ["::imex::xegpu::XeGPUDialect",
+                           "::mlir::vector::VectorDialect",
+                            "::mlir::memref::MemRefDialect",
+                            "::mlir::LLVM::LLVMDialect",
+                            "::mlir::func::FuncDialect",
+                           "::mlir::arith::ArithDialect",
+                            ];
+  let constructor = "imex::createConvertXeGPUToVCPass()";
+}
+
 #endif // _IMEX_CONVERSION_PASSES_TD_INCLUDED_

include/imex/Conversion/XeGPUToVC/CMakeLists.txt

@@ -0,0 +1,44 @@
+//===- XeGPUToVC.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 XeGPU dialect operations into Func
+/// dialect calls to vc-intrinsics functions
+///
+//===----------------------------------------------------------------------===//
+#ifndef IMEX_CONVERSION_XEGPUTOVC_H
+#define IMEX_CONVERSION_XEGPUTOVC_H
+#include <mlir/Dialect/Vector/IR/VectorOps.h>
+
+#include "imex/Dialect/XeGPU/IR/XeGPU.h"
+#include "imex/Dialect/XeTile/IR/XeTileOps.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 {
+
+std::unique_ptr<::mlir::OperationPass<::mlir::gpu::GPUModuleOp>>
+createConvertXeGPUToVCPass();
+
+} // namespace imex
+#endif

include/imex/Conversion/XeGPUToVC/XeGPUToVC.h

@@ -16,13 +16,15 @@
 #ifndef _IMEX_XECOMMON_H_
 #define _IMEX_XECOMMON_H_
 
+#include "imex/Dialect/XeGPU/IR/XeGPU.h"
+#include "imex/Dialect/XeTile/IR/XeTileOps.h"
 #include <mlir/Dialect/GPU/IR/GPUDialect.h>
 #include <mlir/Dialect/SCF/IR/SCF.h>
 #include <mlir/IR/BuiltinOps.h>
 #include <mlir/IR/Value.h>
 #include <mlir/Transforms/DialectConversion.h>
 #include <mlir/Transforms/OneToNTypeConversion.h>
-
+using namespace imex::xegpu;
 namespace imex {
 
 // It checks each GPUFuncOp in the module to see
@@ -85,7 +87,7 @@ class TileUsageAnalysis {
             else if (idx == 2)
               Usage[op] |= (uint)UsageType::DPAS_C;
             else
-              llvm::dbgs() << "unknown usage: " << idx;
+              op->emitOpError() << "unknown usage: " << idx;
           }
 
           if (auto unpack =
@@ -269,6 +271,108 @@ class PropagateAnalysis {
   }
 };
 
+std::pair<std::string, mlir::VectorType>
+encodeVectorType(mlir::ConversionPatternRewriter &rewriter,
+                 mlir::VectorType type, bool use64bitData = false,
+                 bool enforceInteger = false);
+
+unsigned encodeDataum(mlir::Type type);
+
+unsigned encodeOpcode(AtomicRMWKind kind);
+
+// L1 and L3 Cache Policies for Load Operation
+//  L1 Cache Policies: Uncached (UC), Cached (C), Cache Streaming (S),
+//  Invalidate-After-Read (IAR) L3 Cache Policies: Uncached (UC), Cached (C)
+#define L1UC_L3UC 1
+#define L1UC_L3C 2
+#define L1C_L3UC 3
+#define L1C_L3C 4
+#define L1S_L3UC 5
+#define L1S_L3C 6
+#define L1IAR_L3C 7
+
+// L1 and L3 Cache Policies for Store operation
+//  L1 Cache Policies: Uncached (UC), Write-Through (WT), Write-Back (WB),
+//  Streaming (S) L3 Cache Policies: Uncached (UC), Cached (WB)
+#define L1UC_L3WB 2
+#define L1WT_L3UC 3
+#define L1WT_L3WB 4
+#define L1S_L3UC 5
+#define L1S_L3WB 6
+#define L1WB_L3WB 7
+
+template <typename OpType> unsigned encodeCacheHint(OpType op) {
+  auto l1hint = op.getL1Hint();
+  auto l3hint = op.getL3Hint();
+
+  constexpr bool isStore = std::is_same_v<OpType, StoreNDOp> ||
+                           std::is_same_v<OpType, StoreScatterOp>;
+  unsigned cacheHint = L1UC_L3UC;
+
+#define SET_CACHEVALUE(hint, cacheHintVal)                                     \
+  hint.has_value() ? hint.value() : cacheHintVal
+
+  if constexpr (!isStore) {
+
+    auto l1CacheValue = SET_CACHEVALUE(l1hint, CacheReadHint::UNCACHED);
+    auto l3CacheValue = SET_CACHEVALUE(l3hint, CacheReadHint::UNCACHED);
+
+// Setting Cache policy override based on L3 Uncached/Cached value for Load
+// operation
+#define SET_L1L3_CACHEREADHINT(cacheHint, l3CacheValue, uncachedVal,           \
+                               cachedVal)                                      \
+  if (l3CacheValue == CacheReadHint::UNCACHED)                                 \
+    cacheHint = uncachedVal;                                                   \
+  else if (l3CacheValue == CacheReadHint::CACHED)                              \
+    cacheHint = cachedVal;
+
+    switch (l1CacheValue) {
+    case CacheReadHint::UNCACHED:
+      SET_L1L3_CACHEREADHINT(cacheHint, l3CacheValue, L1UC_L3UC, L1UC_L3C);
+      break;
+    case CacheReadHint::CACHED:
+      SET_L1L3_CACHEREADHINT(cacheHint, l3CacheValue, L1C_L3UC, L1C_L3C);
+      break;
+    case CacheReadHint::STREAMING:
+      SET_L1L3_CACHEREADHINT(cacheHint, l3CacheValue, L1S_L3UC, L1S_L3C);
+      break;
+    case CacheReadHint::READ_INVALIDATE:
+      if (l3CacheValue == CacheReadHint::CACHED)
+        cacheHint = L1IAR_L3C;
+      break;
+    }
+
+  } else {
+    auto l1CacheValue = SET_CACHEVALUE(l1hint, CacheWriteHint::UNCACHED);
+    auto l3CacheValue = SET_CACHEVALUE(l3hint, CacheWriteHint::UNCACHED);
+
+// Setting Cache policy override based on L3 Uncached/Write-Back value for Store
+// operation
+#define SET_L1L3_CACHEWRITEHINT(cacheHint, l3CacheValue, uncachedVal,          \
+                                cachedVal)                                     \
+  if (l3CacheValue == CacheWriteHint::UNCACHED)                                \
+    cacheHint = uncachedVal;                                                   \
+  else if (l3CacheValue == CacheWriteHint::WRITE_BACK)                         \
+    cacheHint = cachedVal;
+
+    switch (l1CacheValue) {
+    case CacheWriteHint::UNCACHED:
+      SET_L1L3_CACHEWRITEHINT(cacheHint, l3CacheValue, L1UC_L3UC, L1UC_L3WB);
+      break;
+    case CacheWriteHint::WRITE_THROUGH:
+      SET_L1L3_CACHEWRITEHINT(cacheHint, l3CacheValue, L1WT_L3UC, L1WT_L3WB);
+      break;
+    case CacheWriteHint::STREAMING:
+      SET_L1L3_CACHEWRITEHINT(cacheHint, l3CacheValue, L1S_L3UC, L1S_L3WB);
+      break;
+    case CacheWriteHint::WRITE_BACK:
+      if (l3CacheValue == CacheWriteHint::WRITE_BACK)
+        cacheHint = L1WB_L3WB;
+      break;
+    }
+  }
+  return cacheHint;
+}
 class XeTypeConverter : public mlir::OneToNTypeConverter {
 public:
   // friend class XeConversionPattern;

include/imex/Utils/XeCommon.h

@@ -6,3 +6,4 @@ add_subdirectory(GPUToGPUX)
 add_subdirectory(GPUXToLLVM)
 add_subdirectory(XeGPUToSPIRV)
 add_subdirectory(XeTileToXeGPU)
+add_subdirectory(XeGPUToVC)