[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