[PATCH] D78197: [mlir][spirv] Lower memref with dynamic dimensions to runtime arrays

[Lei Zhang via Phabricator via llvm-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGba49096817b7: [mlir][spirv] Lower memref with dynamic dimensions to runtime arrays (authored by antiagainst).

Changed prior to commit:
  https://reviews.llvm.org/D78197?vs=257685&id=258767#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D78197/new/

https://reviews.llvm.org/D78197

Files:
  mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp
  mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
  mlir/test/Conversion/StandardToSPIRV/std-to-spirv.mlir
  mlir/test/Conversion/StandardToSPIRV/std-types-to-spirv.mlir


Index: mlir/test/Conversion/StandardToSPIRV/std-types-to-spirv.mlir
===================================================================
--- mlir/test/Conversion/StandardToSPIRV/std-types-to-spirv.mlir
+++ mlir/test/Conversion/StandardToSPIRV/std-types-to-spirv.mlir
@@ -486,7 +486,7 @@
 
 // -----
 
-// Check that dynamic shapes are not supported.
+// Dynamic shapes
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [], []>,
@@ -494,13 +494,17 @@
      max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>
 } {
 
+// Check that unranked shapes are not supported.
 // CHECK-LABEL: func @unranked_memref
 // CHECK-SAME: memref<*xi32>
 func @unranked_memref(%arg0: memref<*xi32>) { return }
 
 // CHECK-LABEL: func @dynamic_dim_memref
-// CHECK-SAME: memref<8x?xi32>
-func @dynamic_dim_memref(%arg0: memref<8x?xi32>) { return }
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.rtarray<i32, stride=4> [0]>, StorageBuffer>
+// CHECK-SAME: !spv.ptr<!spv.struct<!spv.rtarray<f32, stride=4> [0]>, StorageBuffer>
+func @dynamic_dim_memref(%arg0: memref<8x?xi32>,
+                         %arg1: memref<?x?xf32>)
+{ return }
 
 } // end module
 
Index: mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp
===================================================================
--- mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp
+++ mlir/lib/Dialect/SPIRV/SPIRVLowering.cpp
@@ -331,10 +331,11 @@
 
 static Optional<Type> convertMemrefType(const spirv::TargetEnv &targetEnv,
                                         MemRefType type) {
-  // TODO(ravishankarm) : Handle dynamic shapes.
-  if (!type.hasStaticShape()) {
+  Optional<spirv::StorageClass> storageClass =
+      SPIRVTypeConverter::getStorageClassForMemorySpace(type.getMemorySpace());
+  if (!storageClass) {
     LLVM_DEBUG(llvm::dbgs()
-               << type << " illegal: dynamic shape unimplemented\n");
+               << type << " illegal: cannot convert memory space\n");
     return llvm::None;
   }
 
@@ -345,27 +346,33 @@
     return llvm::None;
   }
 
+  auto arrayElemType = convertScalarType(targetEnv, scalarType, storageClass);
+  if (!arrayElemType)
+    return llvm::None;
+
   Optional<int64_t> scalarSize = getTypeNumBytes(scalarType);
-  Optional<int64_t> memrefSize = getTypeNumBytes(type);
-  if (!scalarSize || !memrefSize) {
+  if (!scalarSize) {
     LLVM_DEBUG(llvm::dbgs()
-               << type << " illegal: cannot deduce element count\n");
+               << type << " illegal: cannot deduce element size\n");
     return llvm::None;
   }
 
-  auto arrayElemCount = *memrefSize / *scalarSize;
+  if (!type.hasStaticShape()) {
+    auto arrayType = spirv::RuntimeArrayType::get(*arrayElemType, *scalarSize);
+    // Wrap in a struct to satisfy Vulkan interface requirements.
+    auto structType = spirv::StructType::get(arrayType, 0);
+    return spirv::PointerType::get(structType, *storageClass);
+  }
 
-  auto storageClass =
-      SPIRVTypeConverter::getStorageClassForMemorySpace(type.getMemorySpace());
-  if (!storageClass) {
+  Optional<int64_t> memrefSize = getTypeNumBytes(type);
+  if (!memrefSize) {
     LLVM_DEBUG(llvm::dbgs()
-               << type << " illegal: cannot convert memory space\n");
+               << type << " illegal: cannot deduce element count\n");
     return llvm::None;
   }
 
-  auto arrayElemType = convertScalarType(targetEnv, scalarType, storageClass);
-  if (!arrayElemType)
-    return llvm::None;
+  auto arrayElemCount = *memrefSize / *scalarSize;
+
   Optional<int64_t> arrayElemSize = getTypeNumBytes(*arrayElemType);
   if (!arrayElemSize) {
     LLVM_DEBUG(llvm::dbgs()


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D78197.258767.patch
Type: text/x-patch
Size: 3621 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20200420/5f58e1a4/attachment.bin>