[Mlir-commits] [mlir] 756d695 - [mlir][StandardToSPIRV] Add support for lowering index_cast to SPIR-V.

[Hanhan Wang]

Author: Hanhan Wang
Date: 2020-05-11T15:41:25-07:00
New Revision: 756d6959d7ac95969cc0127765bdcbee3804547c

URL: https://github.com/llvm/llvm-project/commit/756d6959d7ac95969cc0127765bdcbee3804547c
DIFF: https://github.com/llvm/llvm-project/commit/756d6959d7ac95969cc0127765bdcbee3804547c.diff

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

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

Added: 
    

Modified: 
    mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp
    mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir

Removed: 
    


################################################################################
diff  --git a/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp b/mlir/lib/Conversion/StandardToSPIRV/ConvertStandardToSPIRV.cpp
index ea006c9b7684..fbe02560008a 100644
--- 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 b6e81a427252..16633664c4ea 100644
--- 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