[mlir][StandardToSPIRV] Add support for lowering index_cast to SPIR-V.

Differential Revision: https://reviews.llvm.org/D79644

diff --git a/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp b/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp
index ea006c9..fbe0256 100644 (file)
--- a/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp
+++ b/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp
@@ -867,6 +867,7 @@ void populateStandardToSPIRVPatterns(MLIRContext *context,
       BoolCmpIOpPattern, ConstantCompositeOpPattern, ConstantScalarOpPattern,
       CmpFOpPattern, CmpIOpPattern, IntLoadOpPattern, LoadOpPattern,
       ReturnOpPattern, SelectOpPattern, IntStoreOpPattern, StoreOpPattern,
+      TypeCastingOpPattern<IndexCastOp, spirv::SConvertOp>,
       TypeCastingOpPattern<SIToFPOp, spirv::ConvertSToFOp>,
       TypeCastingOpPattern<ZeroExtendIOp, spirv::UConvertOp>,
       TypeCastingOpPattern<TruncateIOp, spirv::SConvertOp>,

diff --git a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
index b6e81a4..1663366 100644 (file)
--- a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
+++ b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
@@ -461,6 +461,34 @@ module attributes {
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
+// CHECK-LABEL: index_cast1
+func @index_cast1(%arg0: i16) {
+  // CHECK: spv.SConvert %{{.+}} : i16 to i32
+  %0 = index_cast %arg0 : i16 to index
+  return
+}
+
+// CHECK-LABEL: index_cast2
+func @index_cast2(%arg0: index) {
+  // CHECK: spv.SConvert %{{.+}} : i32 to i16
+  %0 = index_cast %arg0 : index to i16
+  return
+}
+
+// CHECK-LABEL: index_cast3
+func @index_cast3(%arg0: i32) {
+  // CHECK-NOT: spv.SConvert
+  %0 = index_cast %arg0 : i32 to index
+  return
+}
+
+// CHECK-LABEL: index_cast4
+func @index_cast4(%arg0: index) {
+  // CHECK-NOT: spv.SConvert
+  %0 = index_cast %arg0 : index to i32
+  return
+}
+
 // CHECK-LABEL: @fpext1
 func @fpext1(%arg0: f16) -> f64 {
   // CHECK: spv.FConvert %{{.*}} : f16 to f64