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

[hanchung]

Rebases to master.

[hanchung]

I think indices could be negative? so switch to using SConvert.

[antiagainst]

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp

773

I think we need to special treat the case where the src/target bitwidth == SPIRVTypeConverter::getIndexType() by just removing this op? OpSConvert requires the source and target bitwidth to be different. Casting from index to i32 will result in an invalid OpSConvert.

[antiagainst]

In D79644#2027720, @hanchung wrote:

I think indices could be negative? so switch to using SConvert.

Yeah the semantics of std.index_cast requires sign extending : https://mlir.llvm.org/docs/Dialects/Standard/#stdindex_cast-indexcastop

[hanchung]

Add test cases for casting between i32 and index.

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp

773

I think this case is already handled in TypeCastingOpPattern?

auto dstType =
    this->typeConverter.convertType(operation.getResult().getType());
if (dstType == operands.front().getType()) {
  // Due to type conversion, we are seeing the same source and target type.
  // Then we can just erase this operation by forwarding its operand.
  rewriter.replaceOp(operation, operands.front());
} else {
  rewriter.template replaceOpWithNewOp<SPIRVOp>(
      operation, dstType, operands, ArrayRef<NamedAttribute>());
}

The first condition will replace it with operand. Let me add more tests for this case.

[antiagainst]

In D79644#2030301, @hanchung wrote:

Add test cases for casting between i32 and index.

Aha, right. Sorry forgot about the detailed impl there. :)

mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp

patterns.insert<
       UnaryAndBinaryOpPattern<UnsignedDivIOp, spirv::UDivOp>,
       UnaryAndBinaryOpPattern<UnsignedRemIOp, spirv::UModOp>,
       UnaryAndBinaryOpPattern<UnsignedShiftRightOp, spirv::ShiftRightLogicalOp>,
       BitwiseOpPattern<AndOp, spirv::LogicalAndOp, spirv::BitwiseAndOp>,
       BitwiseOpPattern<OrOp, spirv::LogicalOrOp, spirv::BitwiseOrOp>,
       BoolCmpIOpPattern, ConstantCompositeOpPattern, ConstantScalarOpPattern,
       CmpFOpPattern, CmpIOpPattern, IntLoadOpPattern, LoadOpPattern,
       ReturnOpPattern, SelectOpPattern, StoreOpPattern,
+      TypeCastingOpPattern<IndexCastOp, spirv::UConvertOp>,
       TypeCastingOpPattern<SIToFPOp, spirv::ConvertSToFOp>,
       TypeCastingOpPattern<ZeroExtendIOp, spirv::UConvertOp>,
       TypeCastingOpPattern<TruncateIOp, spirv::SConvertOp>,
       TypeCastingOpPattern<FPExtOp, spirv::FConvertOp>,
       TypeCastingOpPattern<FPTruncOp, spirv::FConvertOp>, XOrOpPattern>(
       context, typeConverter);
 }
 } // namespace mlir

mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir

 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Int8, Int16, Int64, Float16, Float64], []>,
     {max_compute_workgroup_invocations = 128 : i32,
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
+// CHECK-LABEL: index_cast1
+func @index_cast1(%arg0: i16) {
+  // CHECK: spv.UConvert %{{.+}} : i16 to i32
+  %0 = index_cast %arg0 : i16 to index
+  return
+}
+
+// CHECK-LABEL: index_cast2
+func @index_cast2(%arg0: index) {
+  // CHECK: spv.UConvert %{{.+}} : i32 to i16
+  %0 = index_cast %arg0 : index to i16
+  return
+}
+
 // CHECK-LABEL: @fpext1
 func @fpext1(%arg0: f16) -> f64 {
   // CHECK: spv.FConvert %{{.*}} : f16 to f64
   %0 = std.fpext %arg0 : f16 to f64
   return %0 : f64
 }
 
 // CHECK-LABEL: @fpext2