[Mlir-commits] [mlir] [MLIR] SM_90 integratation test of TMA with 128b Swizzling (PR #65953)

llvmlistbot at llvm.org llvmlistbot at llvm.org
Mon Sep 11 05:25:22 PDT 2023


llvmbot wrote:

@llvm/pr-subscribers-mlir

<details>
<summary>Changes</summary>

An integration test for the 128b Swizzling TMA.

TMA with 128B Swizzle loads data as follows (each numbered cell is 16 bytes). The program tests this pattern for `128x64xf16` type.

```
|-------------------------------|
| 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 |
| 1 | 0 | 3 | 2 | 5 | 4 | 7 | 6 |
| 2 | 3 | 0 | 1 | 6 | 7 | 4 | 5 |
| 3 | 2 | 1 | 0 | 7 | 6 | 5 | 4 |
| 4 | 5 | 6 | 7 | 0 | 1 | 2 | 3 |
| 5 | 4 | 7 | 6 | 1 | 0 | 3 | 2 |
| 6 | 7 | 4 | 5 | 2 | 3 | 0 | 1 |
|-------------------------------|
| ... pattern repeats ...       |
|-------------------------------|
```
--

Patch is 62.80 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/65953.diff

1 Files Affected:

- (added) mlir/test/Integration/GPU/CUDA/sm90/tmaload_128_64_swizzle128b.mlir (+263) 


<pre>
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tmaload_128_64_swizzle128b.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tmaload_128_64_swizzle128b.mlir
new file mode 100644
index 000000000000000..363a3641f646673
--- /dev/null
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tmaload_128_64_swizzle128b.mlir
@@ -0,0 +1,263 @@
+// RUN: mlir-opt %s --convert-nvgpu-to-nvvm \
+// RUN:         -convert-linalg-to-loops \
+// RUN:         -gpu-kernel-outlining \
+// RUN:         -convert-vector-to-scf  \
+// RUN:         -lower-affine \
+// RUN:         -convert-scf-to-cf \
+// RUN:         -convert-nvvm-to-llvm \
+// RUN:         -convert-nvgpu-to-nvvm \
+// RUN:         -convert-scf-to-cf  \
+// RUN:         -convert-vector-to-llvm \
+// RUN:         -convert-math-to-llvm \
+// RUN:         -convert-index-to-llvm=index-bitwidth=32 \
+// RUN:         -convert-arith-to-llvm \
+// RUN:         -finalize-memref-to-llvm='use-opaque-pointers=1' \
+// RUN:         -convert-func-to-llvm \
+// RUN:         -canonicalize -cse \
+// RUN:         -expand-strided-metadata --nvvm-attach-target="module=main_kernel features=+ptx80 chip=sm_90 O=3" \
+// RUN:  | mlir-opt -pass-pipeline='builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,convert-index-to-llvm{index-bitwidth=32},canonicalize,cse))' \
+// RUN:  | mlir-opt --gpu-to-llvm --gpu-module-to-binary -canonicalize -cse -reconcile-unrealized-casts \
+// RUN:  | mlir-translate --mlir-to-llvmir -o %t.ll 
+// RUN: clang %t.ll -O3 %mlir_cuda_runtime %mlir_runner_utils -o %t.exe
+// RUN: LD_LIBRARY_PATH=%shlibdir %t.exe | FileCheck %s
+
+
+// Test swizzling with TMA load
+// 128B Swizzle Each numbered cell is 16 byte 
+// |-------------------------------|
+// | 0 | 1 | 2 | 3 | 4 | 5 | 6 | 7 | 
+// | 1 | 0 | 3 | 2 | 5 | 4 | 7 | 6 |
+// | 2 | 3 | 0 | 1 | 6 | 7 | 4 | 5 |
+// | 3 | 2 | 1 | 0 | 7 | 6 | 5 | 4 | 
+// | 4 | 5 | 6 | 7 | 0 | 1 | 2 | 3 |
+// | 5 | 4 | 7 | 6 | 1 | 0 | 3 | 2 | 
+// | 6 | 7 | 4 | 5 | 2 | 3 | 0 | 1 |
+// |-------------------------------| 
+// | ... pattern repeats ...       |
+// |-------------------------------|
+
+
+!barrierType = !nvgpu.mbarrier.barrier<memorySpace = #gpu.address_space<workgroup>>
+!tokenType = !nvgpu.mbarrier.token
+
+!lhs = memref<128x64xf16>
+!shmemlhs = memref<128x64xf16, 3>
+!lhsTensorMap = !nvgpu.tensormap.descriptor<tensor = !shmemlhs, swizzle = swizzle_128b, l2promo=none, oob=zero, interleave=none>
+
+module @mymod {
+  func.func private @printMemrefF32(memref<*xf32>)
+  memref.global "private" @bufferLhsGlobal : !shmemlhs
+  llvm.func @printf(!llvm.ptr<i8>, ...) -> i32
+  func.func @main() {
+    %c8192 = arith.constant 8192 : index
+    %c-1_i32 = arith.constant -1 : i32
+    %c10000000 = arith.constant 10000000 : index
+    %c64 = arith.constant 64 : index
+    %c1 = arith.constant 1 : index
+    %c0 = arith.constant 0 : index
+    %c128 = arith.constant 128 : index
+    %c8 = arith.constant 8 : index
+    
+    // Step 1. Allocate host data and initialize it.
+    %alloc = memref.alloc() : !lhs    
+    %alloc_1 = memref.alloc() : memref<128x64xf32>
+    scf.for %arg0 = %c0 to %c128 step %c1 {
+      scf.for %arg1 = %c0 to %c64 step %c1 {
+        %5 = arith.muli %arg0, %c64 : index
+        %6 = arith.addi %5, %arg1 : index
+        %7 = arith.divui %6, %c8 : index
+        %8 = arith.index_cast %7 : index to i32
+        %9 = arith.sitofp %8 : i32 to f16
+        memref.store %9, %alloc[%arg0, %arg1] : !lhs
+        %10 = arith.extf %9 : f16 to f32
+        memref.store %10, %alloc_1[%arg0, %arg1] : memref<128x64xf32>
+      }
+    }
+
+    // Step 2. Print on the host
+    %cast = memref.cast %alloc_1 : memref<128x64xf32> to memref<*xf32>
+    call @printMemrefF32(%cast) : (memref<*xf32>) -> ()
+    
+    // Step 3. Copy host to device
+    %0 = gpu.wait async
+    %memref, %asyncToken = gpu.alloc async [%0] () : !lhs    
+    %1 = gpu.memcpy async [%0] %memref, %alloc : !lhs, !lhs    
+    %cast_6 = memref.cast %memref : !lhs to memref<*xf16>
+    
+    // Step 4. Create TMA tensor descriptor
+    %3 = nvgpu.tma.create.descriptor %cast_6 box[%c128, %c64] : memref<*xf16> -> !lhsTensorMap
+
+    // Step 5. Launch a GPU kernel
+    gpu.launch blocks(%arg0, %arg1, %arg2) in (%arg6 = %c1, %arg7 = %c1, %arg8 = %c1) threads(%arg3, %arg4, %arg5) in (%arg9 = %c128, %arg10 = %c1, %arg11 = %c1) {
+      %5 = gpu.block_dim  x
+      %6 = gpu.thread_id  x
+      %7 = memref.get_global @bufferLhsGlobal : !shmemlhs
+    
+      // Step 6. Initialize the mbarrier
+      %9 = nvgpu.mbarrier.create -> !barrierType
+      nvgpu.mbarrier.init %9, %5 : !barrierType
+      %10 = arith.cmpi eq, %6, %c0 : index
+      
+      
+      // Step 7. First thread does TMA load
+      scf.if %10 {
+        gpu.printf "[GPU] TMA SIZE %d\0A" %c8192 : index
+        nvgpu.tma.async.load %3[%c0, %c0], %9 to %7 : !lhsTensorMap, !barrierType -> !shmemlhs
+        nvgpu.mbarrier.arrive.expect_tx %9, %c8192 : !barrierType
+      } else {
+        nvgpu.mbarrier.arrive.expect_tx %9, %c0 : !barrierType
+      }
+
+      // Step 8. Wait until TMA is done
+      nvgpu.mbarrier.try_wait.parity %9, %c0, %c10000000 : !barrierType
+
+      // Step 9. Print loaded data in 128b swizzled
+      scf.if %10 {        
+        gpu.printf "===--- Matrix A ---=== %d \0A" %c-1_i32 : i32
+        scf.for %arg12 = %c0 to %c128 step %c1 {
+          scf.for %arg13 = %c0 to %c64 step %c1 {
+            %15 = memref.load %7[%arg12, %arg13] : !shmemlhs
+            %16 = arith.extf %15 : f16 to f32
+            gpu.printf "%.0f,   " %16 : f32
+          }
+          gpu.printf "%d\0A" %c-1_i32 : i32
+        }
+        gpu.printf "===----------------=== %d \0A" %c-1_i32 : i32
+      }
+      gpu.terminator
+    }
+    return
+  }
+}
+
+
+// CHECK: [GPU] TMA SIZE 8192
+// CHECK: ===--- Matrix A ---=== -1 
+// CHECK: 0,   0,   0,   0,   0,   0,   0,   0,   1,   1,   1,   1,   1,   1,   1,   1,   2,   2,   2,   2,   2,   2,   2,   2,   3,   3,   3,   3,   3,   3,   3,   3,   4,   4,   4,   4,   4,   4,   4,   4,   5,   5,   5,   5,   5,   5,   5,   5,   6,   6,   6,   6,   6,   6,   6,   6,   7,   7,   7,   7,   7,   7,   7,   7,   -1
+// CHECK: 9,   9,   9,   9,   9,   9,   9,   9,   8,   8,   8,   8,   8,   8,   8,   8,   11,   11,   11,   11,   11,   11,   11,   11,   10,   10,   10,   10,   10,   10,   10,   10,   13,   13,   13,   13,   13,   13,   13,   13,   12,   12,   12,   12,   12,   12,   12,   12,   15,   15,   15,   15,   15,   15,   15,   15,   14,   14,   14,   14,   14,   14,   14,   14,   -1
+// CHECK: 18,   18,   18,   18,   18,   18,   18,   18,   19,   19,   19,   19,   19,   19,   19,   19,   16,   16,   16,   16,   16,   16,   16,   16,   17,   17,   17,   17,   17,   17,   17,   17,   22,   22,   22,   22,   22,   22,   22,   22,   23,   23,   23,   23,   23,   23,   23,   23,   20,   20,   20,   20,   20,   20,   20,   20,   21,   21,   21,   21,   21,   21,   21,   21,   -1
+// CHECK: 27,   27,   27,   27,   27,   27,   27,   27,   26,   26,   26,   26,   26,   26,   26,   26,   25,   25,   25,   25,   25,   25,   25,   25,   24,   24,   24,   24,   24,   24,   24,   24,   31,   31,   31,   31,   31,   31,   31,   31,   30,   30,   30,   30,   30,   30,   30,   30,   29,   29,   29,   29,   29,   29,   29,   29,   28,   28,   28,   28,   28,   28,   28,   28,   -1
+// CHECK: 36,   36,   36,   36,   36,   36,   36,   36,   37,   37,   37,   37,   37,   37,   37,   37,   38,   38,   38,   38,   38,   38,   38,   38,   39,   39,   39,   39,   39,   39,   39,   39,   32,   32,   32,   32,   32,   32,   32,   32,   33,   33,   33,   33,   33,   33,   33,   33,   34,   34,   34,   34,   34,   34,   34,   34,   35,   35,   35,   35,   35,   35,   35,   35,   -1
+// CHECK: 45,   45,   45,   45,   45,   45,   45,   45,   44,   44,   44,   44,   44,   44,   44,   44,   47,   47,   47,   47,   47,   47,   47,   47,   46,   46,   46,   46,   46,   46,   46,   46,   41,   41,   41,   41,   41,   41,   41,   41,   40,   40,   40,   40,   40,   40,   40,   40,   43,   43,   43,   43,   43,   43,   43,   43,   42,   42,   42,   42,   42,   42,   42,   42,   -1
+// CHECK: 54,   54,   54,   54,   54,   54,   54,   54,   55,   55,   55,   55,   55,   55,   55,   55,   52,   52,   52,   52,   52,   52,   52,   52,   53,   53,   53,   53,   53,   53,   53,   53,   50,   50,   50,   50,   50,   50,   50,   50,   51,   51,   51,   51,   51,   51,   51,   51,   48,   48,   48,   48,   48,   48,   48,   48,   49,   49,   49,   49,   49,   49,   49,   49,   -1
+// CHECK: 63,   63,   63,   63,   63,   63,   63,   63,   62,   62,   62,   62,   62,   62,   62,   62,   61,   61,   61,   61,   61,   61,   61,   61,   60,   60,   60,   60,   60,   60,   60,   60,   59,   59,   59,   59,   59,   59,   59,   59,   58,   58,   58,   58,   58,   58,   58,   58,   57,   57,   57,   57,   57,   57,   57,   57,   56,   56,   56,   56,   56,   56,   56,   56,   -1
+// CHECK: 64,   64,   64,   64,   64,   64,   64,   64,   65,   65,   65,   65,   65,   65,   65,   65,   66,   66,   66,   66,   66,   66,   66,   66,   67,   67,   67,   67,   67,   67,   67,   67,   68,   68,   68,   68,   68,   68,   68,   68,   69,   69,   69,   69,   69,   69,   69,   69,   70,   70,   70,   70,   70,   70,   70,   70,   71,   71,   71,   71,   71,   71,   71,   71,   -1
+// CHECK: 73,   73,   73,   73,   73,   73,   73,   73,   72,   72,   72,   72,   72,   72,   72,   72,   75,   75,   75,   75,   75,   75,   75,   75,   74,   74,   74,   74,   74,   74,   74,   74,   77,   77,   77,   77,   77,   77,   77,   77,   76,   76,   76,   76,   76,   76,   76,   76,   79,   79,   79,   79,   79,   79,   79,   79,   78,   78,   78,   78,   78,   78,   78,   78,   -1
+// CHECK: 82,   82,   82,   82,   82,   82,   82,   82,   83,   83,   83,   83,   83,   83,   83,   83,   80,   80,   80,   80,   80,   80,   80,   80,   81,   81,   81,   81,   81,   81,   81,   81,   86,   86,   86,   86,   86,   86,   86,   86,   87,   87,   87,   87,   87,   87,   87,   87,   84,   84,   84,   84,   84,   84,   84,   84,   85,   85,   85,   85,   85,   85,   85,   85,   -1
+// CHECK: 91,   91,   91,   91,   91,   91,   91,   91,   90,   90,   90,   90,   90,   90,   90,   90,   89,   89,   89,   89,   89,   89,   89,   89,   88,   88,   88,   88,   88,   88,   88,   88,   95,   95,   95,   95,   95,   95,   95,   95,   94,   94,   94,   94,   94,   94,   94,   94,   93,   93,   93,   93,   93,   93,   93,   93,   92,   92,   92,   92,   92,   92,   92,   92,   -1
+// CHECK: 100,   100,   100,   100,   100,   100,   100,   100,   101,   101,   101,   101,   101,   101,   101,   101,   102,   102,   102,   102,   102,   102,   102,   102,   103,   103,   103,   103,   103,   103,   103,   103,   96,   96,   96,   96,   96,   96,   96,   96,   97,   97,   97,   97,   97,   97,   97,   97,   98,   98,   98,   98,   98,   98,   98,   98,   99,   99,   99,   99,   99,   99,   99,   99,   -1
+// CHECK: 109,   109,   109,   109,   109,   109,   109,   109,   108,   108,   108,   108,   108,   108,   108,   108,   111,   111,   111,   111,   111,   111,   111,   111,   110,   110,   110,   110,   110,   110,   110,   110,   105,   105,   105,   105,   105,   105,   105,   105,   104,   104,   104,   104,   104,   104,   104,   104,   107,   107,   107,   107,   107,   107,   107,   107,   106,   106,   106,   106,   106,   106,   106,   106,   -1
+// CHECK: 118,   118,   118,   118,   118,   118,   118,   118,   119,   119,   119,   119,   119,   119,   119,   119,   116,   116,   116,   116,   116,   116,   116,   116,   117,   117,   117,   117,   117,   117,   117,   117,   114,   114,   114,   114,   114,   114,   114,   114,   115,   115,   115,   115,   115,   115,   115,   115,   112,   112,   112,   112,   112,   112,   112,   112,   113,   113,   113,   113,   113,   113,   113,   113,   -1
+// CHECK: 127,   127,   127,   127,   127,   127,   127,   127,   126,   126,   126,   126,   126,   126,   126,   126,   125,   125,   125,   125,   125,   125,   125,   125,   124,   124,   124,   124,   124,   124,   124,   124,   123,   123,   123,   123,   123,   123,   123,   123,   122,   122,   122,   122,   122,   122,   122,   122,   121,   121,   121,   121,   121,   121,   121,   121,   120,   120,   120,   120,   120,   120,   120,   120,   -1
+// CHECK: 128,   128,   128,   128,   128,   128,   128,   128,   129,   129,   129,   129,   129,   129,   129,   129,   130,   130,   130,   130,   130,   130,   130,   130,   131,   131,   131,   131,   131,   131,   131,   131,   132,   132,   132,   132,   132,   132,   132,   132,   133,   133,   133,   133,   133,   133,   133,   133,   134,   134,   134,   134,   134,   134,   134,   134,   135,   135,   135,   135,   135,   135,   135,   135,   -1
+// CHECK: 137,   137,   137,   137,   137,   137,   137,   137,   136,   136,   136,   136,   136,   136,   136,   136,   139,   139,   139,   139,   139,   139,   139,   139,   138,   138,   138,   138,   138,   138,   138,   138,   141,   141,   141,   141,   141,   141,   141,   141,   140,   140,   140,   140,   140,   140,   140,   140,   143,   143,   143,   143,   143,   143,   143,   143,   142,   142,   142,   142,   142,   142,   142,   142,   -1
+// CHECK: 146,   146,   146,   146,   146,   146,   146,   146,   147,   147,   147,   147,   147,   147,   147,   147,   144,   144,   144,   144,   144,   144,   144,   144,   145,   145,   145,   145,   145,   145,   145,   145,   150,   150,   150,   150,   150,   150,   150,   150,   151,   151,   151,   151,   151,   151,   151,   151,   148,   148,   148,   148,   148,   148,   148,   148,   149,   149,   149,   149,   149,   149,   149,   149,   -1
+// CHECK: 155,   155,   155,   155,   155,   155,   155,   155,   154,   154,   154,   154,   154,   154,   154,   154,   153,   153,   153,   153,   153,   153,   153,   153,   152,   152,   152,   152,   152,   152,   152,   152,   159,   159,   159,   159,   159,   159,   159,   159,   158,   158,   158,   158,   158,   158,   158,   158,   157,   157,   157,   157,   157,   157,   157,   157,   156,   156,   156,   156,   156,   156,   156,   156,   -1
+// CHECK: 164,   164,   164,   164,   164,   164,   164,   164,   165,   165,   165,   165,   165,   165,   165,   165,   166,   166,   166,   166,   166,   166,   166,   166,   167,   167,   167,   167,   167,   167,   167,   167,   160,   160,   160,   160,   160,   160,   160,   160,   161,   161,   161,   161,   161,   161,   161,   161,   162,   162,   162,   162,   162,   162,   162,   162,   163,   163,   163,   163,   163,   163,   163,   163,   -1
+// CHECK: 173,   173,   173,   173,   173,   173,   173,   173,   172,   172,   172,   172,   172,   172,   172,   172,   175,   175,   175,   175,   175,   175,   175,   175,   174,   174,   174,   174,   174,   174,   174,   174,   169,   169,   169,   169,   169,   169,   169,   169,   168,   168,   168,   168,   168,   168,   168,   168,   171,   171,   171,   171,   171,   171,   171,   171,   170,   170,   170,   170,   170,   170,   170,   170,   -1
+// CHECK: 182,   182,   182,   182,   182,   182,   182,   182,   183,   183,   183,   183,   183,   183,   183,   183,   180,   180,   180,   180,   180,   180,   180,   180,   181,   181,   181,   181,   181,   181,   181,   181,   178,   178,   178,   178,   178,   178,   178,   178,   179,   179,   179,   179,   179,   179,   179,   179,   176,   176,   176,   176,   176,   176,   176,   176,   177,   177,   177,   177,   177,   177,   177,   177,   -1
+// CHECK: 191,   191,   191,   191,   191,   191,   191,   191,   190,   190,   190,   190,   190,   190,   190,   190,   189,   189,   189,   189,   189,   189,   189,   189,   188,   188,   188,   188,   188,   188,   188,   188,   187,   187,   187,   187,   187,   187,   187,   187,   186,   186,   186,   186,   186,   186,   186,   186,   185,   185,   185,   185,   185,   185,   185,   185,   184,   184,   184,   184,   184,   184,   184,   184,   -1
+// CHECK: 192,   192,   192,   192,   192,   192,   192,   192,   193,   193,   193,   193,   193,   193,   193,   193,   194,   194,   194,   194,   194,   194,   194,   194,   195,   195,   195,   195,   195,   195,   195,   195,   196,   196,   196,   196,   196,   196,   196,   196,   197,   197,   197,   197,   197,   197,   197,   197,   198,   198,   198,   198,   198,   198,   198,   198,   199,   199,   199,   199,   199,   199,   199,   199,   -1
+// CHECK: 201,   201,   201,   201,   201,   201,   201,   201,   200,   200,   200,   200,   200,   200,   200,   200,   203,   203,   203,   203,   203,   203,   203,   203,   202,   202,   202,   202,   202,   202,   202,   202,   205,   205,   205,   205,   205,   205,   205,   205,   204,   204,   204,   204,   204,   204,   204,   204,   207,   207,   207,   207,   207,   207,   207,   207,   206,   206,   206,   206,   206,   206,   206,   206,   -1
+// CHECK: 210,   210,   210,   210,   210,   210,   210,   210,   211,   211,   211,   211,   211,   211,   211,   211,   208,   208,   208,   208,   208,   208,   208,   208,   209,   209,   209,   209,   209,   209,   209,   209,   214,   214,   214,   214,   214,   214,   214,   214,   215,   215,   215,   215,   215,   215,   215,   215,   212,   212,   212,   212,   212,   212,   212,   212,   213,   213,   213,   213,   213,   213,   213,   213,   -1
+// CHECK: 219,   219,   219,   219,   219,   219,   219,   219,   218,   218,   218,   218,   218,   218,   218,   218,   217,   217,   217,   217,   217,   217,   217,   217,   216,   216,   216,   216,   216,   216,   216,   216,   223,   223,   223,   223,   223,   223,   223,   223,   222,   222,   222,   222,   222,   222,   222,   222,   221,   221,   221,   221,   221,   221,   221,   221,   220,   220,   220,   220,   220,   220,   220,   220,   -1
+// CHECK: 228,   228,   228,   228,   228,   228,   228,   228,   229,   229,   229,   229,   229,   229,   229,   229,   230,   230,   230,   230,   230,   230,   230,   230,   231,   231,   231,   231,   231,   231,   231,   231,   224,   224,   224,   224,   224,   224,   224,   224,   225,   225,   225,   225,   225,   225,   225,   225,   226,   226,   226,   226,   226,   226,   226,   226,   227,   227,   227,   227,   227,   227,   227,   227,   -1
+// CHECK: 237,   237,   237,   237,   237,   237,   237,   237,   236,   236,   236,   236,   236,   236,   236,   236,   239,   239,   239,   239,   239,   239,   239,   239,   238,   238,   238,   238,   238,   238,   238,   238,   233,   233,   233,   233,   233,   233,   233,   233,   232,   232,   232,   232,   232,   232,   232,   232,   235,   235,   235,   235,   235,   235,   235,   235,   234,   234,   234,   234,   234,   234,   234,   234,   -1
+// CHECK: 246,   246,   246,   246,   246,   246,   246,   246,   247,   247,   247,   247,   247,   247,   247,   247,   244,   244,   244,   244,   244,   244,   244,   244,   245,   245,   245,   245,   245,   245,   245,   245,   242,   242,   242,   242,   242,   242,   242,   242,   243,   243,   243,   243,   243,   243,   243,   243,   240,   240,   240,   240,   240,   240,   240,   240,   241,   241,   241,   241,   241,   241,   241,   241,   -1
+// CHECK: 255,   255,   255,   255,   255,   255,   255,   255,   254,   254,   254,   254,   254,   254,   254,   254,   253,   253,   253,   253,   253,   253,   253,   253,   252,   252,   252,   252,   252,   252,   252,   252,   251,   251,   251,   251,   251,   251,   251,   251,   250,   250,   250,   250,   250,   250,   250,   250,   249,   249,   249,   249,   249,   249,   249,   249,   248,   248,   248,   248,   248,   248,   248,   248,   -1
+// CHECK: 256,   256,   256,   256,   256,   256,   256,   256,   257,   257,   257,   257,   257,   257,   257,   257,   258,   258,   258,   258,   258,   258,   258,   258,   259,   259,   259,   259,   259,   259,   259,   259,   260,   260,   260,   260,   260,   260,   260,   260,   261,   261,   261,   261,   261,   261,   261,   261,   262,   262,   262,   262,   262,   262,   262,   262,   263,   263,   263,   263,   263,   263,   263,   263,   -1
+// CHECK: 265,   265,   265,   265,   265,   265,   265,   265,   264,   264,   264,   264,   264,   264,   264,   264,   267,   267,   267,   267,   267,   267,   26...
<truncated>
</pre>

</details>

https://github.com/llvm/llvm-project/pull/65953


More information about the Mlir-commits mailing list