[Mlir-commits] [mlir] [MLIR] SM_90 integratation test of TMA `128x64xf16` and `64x64xf16` with 128b Swizzling (PR #65954)
Guray Ozen
llvmlistbot at llvm.org
Fri Sep 15 00:30:07 PDT 2023
https://github.com/grypp updated https://github.com/llvm/llvm-project/pull/65954
>From 6f8184377299671df2ff30072c5fa0eb6f97b181 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.ozen at gmail.com>
Date: Mon, 11 Sep 2023 14:29:04 +0200
Subject: [PATCH 1/2] [MLIR] SM_90 integratation test of TMA `128x64xf16` and
`64x64xf16` with 128b Swizzling
The #65953 added a test `128x64xf16` that does a single TMA load. This PR adds more complex test that does 2 additional TMA loads with 128B Swizzling:
```
TMA Load: Matrix-A[0:128][0:64]
TMA Load: Matrix-B[0:64][0:64]
TMA Load: Matrix-B[64:128][0:64]
```
The program tests the loaded data for Matrix-B.
---
.../CUDA/sm90/tmaload_64_64_swizzle128b.mlir | 245 ++++++++++++++++++
1 file changed, 245 insertions(+)
create mode 100644 mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir
new file mode 100644
index 000000000000000..b3c55c87405c507
--- /dev/null
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir
@@ -0,0 +1,245 @@
+// RUN: mlir-opt %s --convert-nvgpu-to-nvvm \
+// RUN: -convert-linalg-to-loops \
+// RUN: -canonicalize -cse \
+// RUN: -gpu-kernel-outlining \
+// RUN: -canonicalize -cse \
+// RUN: -convert-vector-to-scf \
+// RUN: -canonicalize -cse \
+// RUN: -lower-affine \
+// RUN: -canonicalize -cse \
+// RUN: -convert-scf-to-cf \
+// RUN: -canonicalize -cse \
+// RUN: -convert-nvvm-to-llvm \
+// RUN: -canonicalize -cse \
+// RUN: -convert-nvgpu-to-nvvm \
+// RUN: -canonicalize -cse \
+// RUN: -convert-scf-to-cf \
+// RUN: -convert-vector-to-llvm \
+// RUN: -canonicalize -cse \
+// RUN: -convert-math-to-llvm \
+// RUN: -canonicalize -cse \
+// RUN: -lower-affine \
+// 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
+
+// This does 3 TMA loads with 128B Swizzling :
+// TMA Load: Matrix-A[0:128][0:64]
+// TMA Load: Matrix-B[0:64][0:64]
+// TMA Load: Matrix-B[64:128][0:64]
+
+// 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>
+
+!rhs = memref<128x64xf16>
+!shmemrhs = memref<128x64xf16, 3>
+!rhsTensorMap = !nvgpu.tensormap.descriptor<tensor = !shmemrhs, swizzle = swizzle_128b, l2promo=none, oob=zero, interleave=none>
+
+module @mymod {
+ func.func private @printMemrefF32(memref<*xf32>)
+ memref.global "private" @bufferLhsGlobal : !shmemlhs
+ memref.global "private" @bufferRhsGlobal : !shmemrhs
+ llvm.func @printf(!llvm.ptr<i8>, ...) -> i32
+ func.func @main() {
+ %c32768 = arith.constant 32768 : index
+ %c-1_i32 = arith.constant -1 : i32
+ %c10000000 = arith.constant 10000000 : index
+ %c64 = arith.constant 64 : index
+ %c1 = arith.constant 1 : index
+ %c32 = arith.constant 32 : index
+ %c0 = arith.constant 0 : index
+ %c128 = arith.constant 128 : index
+ %c8 = arith.constant 8 : index
+
+ // Step 1. Allocate host data and initialize it.
+ %lhs = memref.alloc() : !lhs
+ %rhs = memref.alloc() : !rhs
+ %lhs32 = memref.alloc() : memref<128x64xf32>
+ %rhs32 = memref.alloc() : memref<64x128xf32>
+ scf.for %i = %c0 to %c64 step %c1 {
+ scf.for %j = %c0 to %c128 step %c1 {
+ %v0 = arith.muli %i, %c128 : index
+ %v00 = arith.addi %v0, %j : index
+ %v01 = arith.divui %v00, %c8 : index
+ %v2 = arith.index_cast %v01 : index to i32
+ %vR = arith.sitofp %v2 : i32 to f16
+ memref.store %vR, %rhs[%i, %j] : !rhs
+ %vR32 = arith.extf %vR : f16 to f32
+ memref.store %vR32, %rhs32[%i, %j] : memref<64x128xf32>
+ }
+ }
+ scf.for %j = %c0 to %c128 step %c1 {
+ scf.for %i = %c0 to %c64 step %c1 {
+ %b0 = arith.muli %j, %c64 : index
+ %b00 = arith.addi %b0, %i : index
+ %b01 = arith.divui %b00, %c8 : index
+ %v1 = arith.index_cast %b01 : index to i32
+ %vL = arith.sitofp %v1 : i32 to f16
+ memref.store %vL, %lhs[%j, %i] : !lhs
+ %vL32 = arith.extf %vL : f16 to f32
+ memref.store %vL32, %lhs32[%j, %i] : memref<128x64xf32>
+ }
+ }
+
+ // Step 2. Print on the host
+ %lhs32_unranked = memref.cast %lhs32 : memref<128x64xf32> to memref<*xf32>
+ call @printMemrefF32(%lhs32_unranked) : (memref<*xf32>) -> ()
+ %rhs32_unranked = memref.cast %rhs32 : memref<64x128xf32> to memref<*xf32>
+ call @printMemrefF32(%rhs32_unranked) : (memref<*xf32>) -> ()
+
+ // Step 3. Copy host to device
+ %0 = gpu.wait async
+ %d_glbmem_lhs, %asyncToken = gpu.alloc async [%0] () : !lhs
+ %d_glbmem_rhs, %asyncToken_2 = gpu.alloc async [%0] () : !rhs
+ %1 = gpu.memcpy async [%0] %d_glbmem_lhs, %lhs : !lhs, !lhs
+ %2 = gpu.memcpy async [%0] %d_glbmem_rhs, %rhs : !rhs, !rhs
+
+ // Step 4. Create TMA tensor descriptor
+ %d_lhs_unranked = memref.cast %d_glbmem_lhs :!lhs to memref<*xf16>
+ %d_rhs_unranked = memref.cast %d_glbmem_rhs :!rhs to memref<*xf16>
+
+ %d_lhsTensorMap = nvgpu.tma.create.descriptor %d_lhs_unranked box[%c128, %c64] : memref<*xf16> -> !lhsTensorMap
+ %d_rhsTensorMap = nvgpu.tma.create.descriptor %d_rhs_unranked box[%c64, %c64] : memref<*xf16> -> !rhsTensorMap
+
+ // 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
+ %lhsShmem = memref.get_global @bufferLhsGlobal : !shmemlhs
+ %rhsShmem = memref.get_global @bufferRhsGlobal : !shmemrhs
+ %rhsShmem2 = memref.subview %rhsShmem[%c32, %c0][%c32, %c128][%c1, %c1] : !shmemrhs to memref<?x?xf16, strided<[?, ?], offset: ?>, 3>
+
+ // 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" %c32768 : index
+ nvgpu.tma.async.load %d_lhsTensorMap[%c0, %c0], %9 to %lhsShmem : !lhsTensorMap, !barrierType -> !shmemlhs
+ nvgpu.tma.async.load %d_rhsTensorMap[%c0, %c0], %9 to %rhsShmem : !rhsTensorMap, !barrierType -> !shmemrhs
+ nvgpu.tma.async.load %d_rhsTensorMap[%c64, %c0], %9 to %rhsShmem2 : !rhsTensorMap, !barrierType -> memref<?x?xf16, strided<[?, ?], offset: ?>, 3>
+ nvgpu.mbarrier.arrive.expect_tx %9, %c32768 : !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 B ---=== %d \n" %c-1_i32 : i32
+ scf.for %ii = %c0 to %c64 step %c1 {
+ scf.for %j = %c0 to %c128 step %c1 {
+ %lhs0 = memref.load %rhsShmem[%ii, %j] : !shmemrhs
+ %lhs032 = arith.extf %lhs0: f16 to f32
+ gpu.printf "%.0f, " %lhs032 : f32
+ }
+ gpu.printf "%d\n" %c-1_i32 : i32
+ }
+ gpu.printf "===----------------=== %d \n" %c-1_i32 : i32
+ }
+ gpu.terminator
+ }
+ return
+ }
+}
+
+
+// CHECK: [GPU] TMA SIZE 32768
+// CHECK: ===--- Matrix B ---=== -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, 17, 17, 17, 17, 17, 17, 17, 17, 16, 16, 16, 16, 16, 16, 16, 16, 19, 19, 19, 19, 19, 19, 19, 19, 18, 18, 18, 18, 18, 18, 18, 18, 21, 21, 21, 21, 21, 21, 21, 21, 20, 20, 20, 20, 20, 20, 20, 20, 23, 23, 23, 23, 23, 23, 23, 23, 22, 22, 22, 22, 22, 22, 22, 22, -1
+// CHECK: 17, 17, 17, 17, 17, 17, 17, 17, 16, 16, 16, 16, 16, 16, 16, 16, 19, 19, 19, 19, 19, 19, 19, 19, 18, 18, 18, 18, 18, 18, 18, 18, 21, 21, 21, 21, 21, 21, 21, 21, 20, 20, 20, 20, 20, 20, 20, 20, 23, 23, 23, 23, 23, 23, 23, 23, 22, 22, 22, 22, 22, 22, 22, 22, 34, 34, 34, 34, 34, 34, 34, 34, 35, 35, 35, 35, 35, 35, 35, 35, 32, 32, 32, 32, 32, 32, 32, 32, 33, 33, 33, 33, 33, 33, 33, 33, 38, 38, 38, 38, 38, 38, 38, 38, 39, 39, 39, 39, 39, 39, 39, 39, 36, 36, 36, 36, 36, 36, 36, 36, 37, 37, 37, 37, 37, 37, 37, 37, -1
+// CHECK: 34, 34, 34, 34, 34, 34, 34, 34, 35, 35, 35, 35, 35, 35, 35, 35, 32, 32, 32, 32, 32, 32, 32, 32, 33, 33, 33, 33, 33, 33, 33, 33, 38, 38, 38, 38, 38, 38, 38, 38, 39, 39, 39, 39, 39, 39, 39, 39, 36, 36, 36, 36, 36, 36, 36, 36, 37, 37, 37, 37, 37, 37, 37, 37, 51, 51, 51, 51, 51, 51, 51, 51, 50, 50, 50, 50, 50, 50, 50, 50, 49, 49, 49, 49, 49, 49, 49, 49, 48, 48, 48, 48, 48, 48, 48, 48, 55, 55, 55, 55, 55, 55, 55, 55, 54, 54, 54, 54, 54, 54, 54, 54, 53, 53, 53, 53, 53, 53, 53, 53, 52, 52, 52, 52, 52, 52, 52, 52, -1
+// CHECK: 51, 51, 51, 51, 51, 51, 51, 51, 50, 50, 50, 50, 50, 50, 50, 50, 49, 49, 49, 49, 49, 49, 49, 49, 48, 48, 48, 48, 48, 48, 48, 48, 55, 55, 55, 55, 55, 55, 55, 55, 54, 54, 54, 54, 54, 54, 54, 54, 53, 53, 53, 53, 53, 53, 53, 53, 52, 52, 52, 52, 52, 52, 52, 52, 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, 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, -1
+// CHECK: 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, 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, 85, 85, 85, 85, 85, 85, 85, 85, 84, 84, 84, 84, 84, 84, 84, 84, 87, 87, 87, 87, 87, 87, 87, 87, 86, 86, 86, 86, 86, 86, 86, 86, 81, 81, 81, 81, 81, 81, 81, 81, 80, 80, 80, 80, 80, 80, 80, 80, 83, 83, 83, 83, 83, 83, 83, 83, 82, 82, 82, 82, 82, 82, 82, 82, -1
+// CHECK: 85, 85, 85, 85, 85, 85, 85, 85, 84, 84, 84, 84, 84, 84, 84, 84, 87, 87, 87, 87, 87, 87, 87, 87, 86, 86, 86, 86, 86, 86, 86, 86, 81, 81, 81, 81, 81, 81, 81, 81, 80, 80, 80, 80, 80, 80, 80, 80, 83, 83, 83, 83, 83, 83, 83, 83, 82, 82, 82, 82, 82, 82, 82, 82, 102, 102, 102, 102, 102, 102, 102, 102, 103, 103, 103, 103, 103, 103, 103, 103, 100, 100, 100, 100, 100, 100, 100, 100, 101, 101, 101, 101, 101, 101, 101, 101, 98, 98, 98, 98, 98, 98, 98, 98, 99, 99, 99, 99, 99, 99, 99, 99, 96, 96, 96, 96, 96, 96, 96, 96, 97, 97, 97, 97, 97, 97, 97, 97, -1
+// CHECK: 102, 102, 102, 102, 102, 102, 102, 102, 103, 103, 103, 103, 103, 103, 103, 103, 100, 100, 100, 100, 100, 100, 100, 100, 101, 101, 101, 101, 101, 101, 101, 101, 98, 98, 98, 98, 98, 98, 98, 98, 99, 99, 99, 99, 99, 99, 99, 99, 96, 96, 96, 96, 96, 96, 96, 96, 97, 97, 97, 97, 97, 97, 97, 97, 119, 119, 119, 119, 119, 119, 119, 119, 118, 118, 118, 118, 118, 118, 118, 118, 117, 117, 117, 117, 117, 117, 117, 117, 116, 116, 116, 116, 116, 116, 116, 116, 115, 115, 115, 115, 115, 115, 115, 115, 114, 114, 114, 114, 114, 114, 114, 114, 113, 113, 113, 113, 113, 113, 113, 113, 112, 112, 112, 112, 112, 112, 112, 112, -1
+// CHECK: 119, 119, 119, 119, 119, 119, 119, 119, 118, 118, 118, 118, 118, 118, 118, 118, 117, 117, 117, 117, 117, 117, 117, 117, 116, 116, 116, 116, 116, 116, 116, 116, 115, 115, 115, 115, 115, 115, 115, 115, 114, 114, 114, 114, 114, 114, 114, 114, 113, 113, 113, 113, 113, 113, 113, 113, 112, 112, 112, 112, 112, 112, 112, 112, 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: 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, 145, 145, 145, 145, 145, 145, 145, 145, 144, 144, 144, 144, 144, 144, 144, 144, 147, 147, 147, 147, 147, 147, 147, 147, 146, 146, 146, 146, 146, 146, 146, 146, 149, 149, 149, 149, 149, 149, 149, 149, 148, 148, 148, 148, 148, 148, 148, 148, 151, 151, 151, 151, 151, 151, 151, 151, 150, 150, 150, 150, 150, 150, 150, 150, -1
+// CHECK: 145, 145, 145, 145, 145, 145, 145, 145, 144, 144, 144, 144, 144, 144, 144, 144, 147, 147, 147, 147, 147, 147, 147, 147, 146, 146, 146, 146, 146, 146, 146, 146, 149, 149, 149, 149, 149, 149, 149, 149, 148, 148, 148, 148, 148, 148, 148, 148, 151, 151, 151, 151, 151, 151, 151, 151, 150, 150, 150, 150, 150, 150, 150, 150, 162, 162, 162, 162, 162, 162, 162, 162, 163, 163, 163, 163, 163, 163, 163, 163, 160, 160, 160, 160, 160, 160, 160, 160, 161, 161, 161, 161, 161, 161, 161, 161, 166, 166, 166, 166, 166, 166, 166, 166, 167, 167, 167, 167, 167, 167, 167, 167, 164, 164, 164, 164, 164, 164, 164, 164, 165, 165, 165, 165, 165, 165, 165, 165, -1
+// CHECK: 162, 162, 162, 162, 162, 162, 162, 162, 163, 163, 163, 163, 163, 163, 163, 163, 160, 160, 160, 160, 160, 160, 160, 160, 161, 161, 161, 161, 161, 161, 161, 161, 166, 166, 166, 166, 166, 166, 166, 166, 167, 167, 167, 167, 167, 167, 167, 167, 164, 164, 164, 164, 164, 164, 164, 164, 165, 165, 165, 165, 165, 165, 165, 165, 179, 179, 179, 179, 179, 179, 179, 179, 178, 178, 178, 178, 178, 178, 178, 178, 177, 177, 177, 177, 177, 177, 177, 177, 176, 176, 176, 176, 176, 176, 176, 176, 183, 183, 183, 183, 183, 183, 183, 183, 182, 182, 182, 182, 182, 182, 182, 182, 181, 181, 181, 181, 181, 181, 181, 181, 180, 180, 180, 180, 180, 180, 180, 180, -1
+// CHECK: 179, 179, 179, 179, 179, 179, 179, 179, 178, 178, 178, 178, 178, 178, 178, 178, 177, 177, 177, 177, 177, 177, 177, 177, 176, 176, 176, 176, 176, 176, 176, 176, 183, 183, 183, 183, 183, 183, 183, 183, 182, 182, 182, 182, 182, 182, 182, 182, 181, 181, 181, 181, 181, 181, 181, 181, 180, 180, 180, 180, 180, 180, 180, 180, 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, 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, -1
+// CHECK: 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, 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, 213, 213, 213, 213, 213, 213, 213, 213, 212, 212, 212, 212, 212, 212, 212, 212, 215, 215, 215, 215, 215, 215, 215, 215, 214, 214, 214, 214, 214, 214, 214, 214, 209, 209, 209, 209, 209, 209, 209, 209, 208, 208, 208, 208, 208, 208, 208, 208, 211, 211, 211, 211, 211, 211, 211, 211, 210, 210, 210, 210, 210, 210, 210, 210, -1
+// CHECK: 213, 213, 213, 213, 213, 213, 213, 213, 212, 212, 212, 212, 212, 212, 212, 212, 215, 215, 215, 215, 215, 215, 215, 215, 214, 214, 214, 214, 214, 214, 214, 214, 209, 209, 209, 209, 209, 209, 209, 209, 208, 208, 208, 208, 208, 208, 208, 208, 211, 211, 211, 211, 211, 211, 211, 211, 210, 210, 210, 210, 210, 210, 210, 210, 230, 230, 230, 230, 230, 230, 230, 230, 231, 231, 231, 231, 231, 231, 231, 231, 228, 228, 228, 228, 228, 228, 228, 228, 229, 229, 229, 229, 229, 229, 229, 229, 226, 226, 226, 226, 226, 226, 226, 226, 227, 227, 227, 227, 227, 227, 227, 227, 224, 224, 224, 224, 224, 224, 224, 224, 225, 225, 225, 225, 225, 225, 225, 225, -1
+// CHECK: 230, 230, 230, 230, 230, 230, 230, 230, 231, 231, 231, 231, 231, 231, 231, 231, 228, 228, 228, 228, 228, 228, 228, 228, 229, 229, 229, 229, 229, 229, 229, 229, 226, 226, 226, 226, 226, 226, 226, 226, 227, 227, 227, 227, 227, 227, 227, 227, 224, 224, 224, 224, 224, 224, 224, 224, 225, 225, 225, 225, 225, 225, 225, 225, 247, 247, 247, 247, 247, 247, 247, 247, 246, 246, 246, 246, 246, 246, 246, 246, 245, 245, 245, 245, 245, 245, 245, 245, 244, 244, 244, 244, 244, 244, 244, 244, 243, 243, 243, 243, 243, 243, 243, 243, 242, 242, 242, 242, 242, 242, 242, 242, 241, 241, 241, 241, 241, 241, 241, 241, 240, 240, 240, 240, 240, 240, 240, 240, -1
+// CHECK: 247, 247, 247, 247, 247, 247, 247, 247, 246, 246, 246, 246, 246, 246, 246, 246, 245, 245, 245, 245, 245, 245, 245, 245, 244, 244, 244, 244, 244, 244, 244, 244, 243, 243, 243, 243, 243, 243, 243, 243, 242, 242, 242, 242, 242, 242, 242, 242, 241, 241, 241, 241, 241, 241, 241, 241, 240, 240, 240, 240, 240, 240, 240, 240, 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: 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, 273, 273, 273, 273, 273, 273, 273, 273, 272, 272, 272, 272, 272, 272, 272, 272, 275, 275, 275, 275, 275, 275, 275, 275, 274, 274, 274, 274, 274, 274, 274, 274, 277, 277, 277, 277, 277, 277, 277, 277, 276, 276, 276, 276, 276, 276, 276, 276, 279, 279, 279, 279, 279, 279, 279, 279, 278, 278, 278, 278, 278, 278, 278, 278, -1
+// CHECK: 273, 273, 273, 273, 273, 273, 273, 273, 272, 272, 272, 272, 272, 272, 272, 272, 275, 275, 275, 275, 275, 275, 275, 275, 274, 274, 274, 274, 274, 274, 274, 274, 277, 277, 277, 277, 277, 277, 277, 277, 276, 276, 276, 276, 276, 276, 276, 276, 279, 279, 279, 279, 279, 279, 279, 279, 278, 278, 278, 278, 278, 278, 278, 278, 290, 290, 290, 290, 290, 290, 290, 290, 291, 291, 291, 291, 291, 291, 291, 291, 288, 288, 288, 288, 288, 288, 288, 288, 289, 289, 289, 289, 289, 289, 289, 289, 294, 294, 294, 294, 294, 294, 294, 294, 295, 295, 295, 295, 295, 295, 295, 295, 292, 292, 292, 292, 292, 292, 292, 292, 293, 293, 293, 293, 293, 293, 293, 293, -1
+// CHECK: 290, 290, 290, 290, 290, 290, 290, 290, 291, 291, 291, 291, 291, 291, 291, 291, 288, 288, 288, 288, 288, 288, 288, 288, 289, 289, 289, 289, 289, 289, 289, 289, 294, 294, 294, 294, 294, 294, 294, 294, 295, 295, 295, 295, 295, 295, 295, 295, 292, 292, 292, 292, 292, 292, 292, 292, 293, 293, 293, 293, 293, 293, 293, 293, 307, 307, 307, 307, 307, 307, 307, 307, 306, 306, 306, 306, 306, 306, 306, 306, 305, 305, 305, 305, 305, 305, 305, 305, 304, 304, 304, 304, 304, 304, 304, 304, 311, 311, 311, 311, 311, 311, 311, 311, 310, 310, 310, 310, 310, 310, 310, 310, 309, 309, 309, 309, 309, 309, 309, 309, 308, 308, 308, 308, 308, 308, 308, 308, -1
+// CHECK: 307, 307, 307, 307, 307, 307, 307, 307, 306, 306, 306, 306, 306, 306, 306, 306, 305, 305, 305, 305, 305, 305, 305, 305, 304, 304, 304, 304, 304, 304, 304, 304, 311, 311, 311, 311, 311, 311, 311, 311, 310, 310, 310, 310, 310, 310, 310, 310, 309, 309, 309, 309, 309, 309, 309, 309, 308, 308, 308, 308, 308, 308, 308, 308, 324, 324, 324, 324, 324, 324, 324, 324, 325, 325, 325, 325, 325, 325, 325, 325, 326, 326, 326, 326, 326, 326, 326, 326, 327, 327, 327, 327, 327, 327, 327, 327, 320, 320, 320, 320, 320, 320, 320, 320, 321, 321, 321, 321, 321, 321, 321, 321, 322, 322, 322, 322, 322, 322, 322, 322, 323, 323, 323, 323, 323, 323, 323, 323, -1
+// CHECK: 324, 324, 324, 324, 324, 324, 324, 324, 325, 325, 325, 325, 325, 325, 325, 325, 326, 326, 326, 326, 326, 326, 326, 326, 327, 327, 327, 327, 327, 327, 327, 327, 320, 320, 320, 320, 320, 320, 320, 320, 321, 321, 321, 321, 321, 321, 321, 321, 322, 322, 322, 322, 322, 322, 322, 322, 323, 323, 323, 323, 323, 323, 323, 323, 341, 341, 341, 341, 341, 341, 341, 341, 340, 340, 340, 340, 340, 340, 340, 340, 343, 343, 343, 343, 343, 343, 343, 343, 342, 342, 342, 342, 342, 342, 342, 342, 337, 337, 337, 337, 337, 337, 337, 337, 336, 336, 336, 336, 336, 336, 336, 336, 339, 339, 339, 339, 339, 339, 339, 339, 338, 338, 338, 338, 338, 338, 338, 338, -1
+// CHECK: 341, 341, 341, 341, 341, 341, 341, 341, 340, 340, 340, 340, 340, 340, 340, 340, 343, 343, 343, 343, 343, 343, 343, 343, 342, 342, 342, 342, 342, 342, 342, 342, 337, 337, 337, 337, 337, 337, 337, 337, 336, 336, 336, 336, 336, 336, 336, 336, 339, 339, 339, 339, 339, 339, 339, 339, 338, 338, 338, 338, 338, 338, 338, 338, 358, 358, 358, 358, 358, 358, 358, 358, 359, 359, 359, 359, 359, 359, 359, 359, 356, 356, 356, 356, 356, 356, 356, 356, 357, 357, 357, 357, 357, 357, 357, 357, 354, 354, 354, 354, 354, 354, 354, 354, 355, 355, 355, 355, 355, 355, 355, 355, 352, 352, 352, 352, 352, 352, 352, 352, 353, 353, 353, 353, 353, 353, 353, 353, -1
+// CHECK: 358, 358, 358, 358, 358, 358, 358, 358, 359, 359, 359, 359, 359, 359, 359, 359, 356, 356, 356, 356, 356, 356, 356, 356, 357, 357, 357, 357, 357, 357, 357, 357, 354, 354, 354, 354, 354, 354, 354, 354, 355, 355, 355, 355, 355, 355, 355, 355, 352, 352, 352, 352, 352, 352, 352, 352, 353, 353, 353, 353, 353, 353, 353, 353, 375, 375, 375, 375, 375, 375, 375, 375, 374, 374, 374, 374, 374, 374, 374, 374, 373, 373, 373, 373, 373, 373, 373, 373, 372, 372, 372, 372, 372, 372, 372, 372, 371, 371, 371, 371, 371, 371, 371, 371, 370, 370, 370, 370, 370, 370, 370, 370, 369, 369, 369, 369, 369, 369, 369, 369, 368, 368, 368, 368, 368, 368, 368, 368, -1
+// CHECK: 375, 375, 375, 375, 375, 375, 375, 375, 374, 374, 374, 374, 374, 374, 374, 374, 373, 373, 373, 373, 373, 373, 373, 373, 372, 372, 372, 372, 372, 372, 372, 372, 371, 371, 371, 371, 371, 371, 371, 371, 370, 370, 370, 370, 370, 370, 370, 370, 369, 369, 369, 369, 369, 369, 369, 369, 368, 368, 368, 368, 368, 368, 368, 368, 384, 384, 384, 384, 384, 384, 384, 384, 385, 385, 385, 385, 385, 385, 385, 385, 386, 386, 386, 386, 386, 386, 386, 386, 387, 387, 387, 387, 387, 387, 387, 387, 388, 388, 388, 388, 388, 388, 388, 388, 389, 389, 389, 389, 389, 389, 389, 389, 390, 390, 390, 390, 390, 390, 390, 390, 391, 391, 391, 391, 391, 391, 391, 391, -1
+// CHECK: 384, 384, 384, 384, 384, 384, 384, 384, 385, 385, 385, 385, 385, 385, 385, 385, 386, 386, 386, 386, 386, 386, 386, 386, 387, 387, 387, 387, 387, 387, 387, 387, 388, 388, 388, 388, 388, 388, 388, 388, 389, 389, 389, 389, 389, 389, 389, 389, 390, 390, 390, 390, 390, 390, 390, 390, 391, 391, 391, 391, 391, 391, 391, 391, 401, 401, 401, 401, 401, 401, 401, 401, 400, 400, 400, 400, 400, 400, 400, 400, 403, 403, 403, 403, 403, 403, 403, 403, 402, 402, 402, 402, 402, 402, 402, 402, 405, 405, 405, 405, 405, 405, 405, 405, 404, 404, 404, 404, 404, 404, 404, 404, 407, 407, 407, 407, 407, 407, 407, 407, 406, 406, 406, 406, 406, 406, 406, 406, -1
+// CHECK: 401, 401, 401, 401, 401, 401, 401, 401, 400, 400, 400, 400, 400, 400, 400, 400, 403, 403, 403, 403, 403, 403, 403, 403, 402, 402, 402, 402, 402, 402, 402, 402, 405, 405, 405, 405, 405, 405, 405, 405, 404, 404, 404, 404, 404, 404, 404, 404, 407, 407, 407, 407, 407, 407, 407, 407, 406, 406, 406, 406, 406, 406, 406, 406, 418, 418, 418, 418, 418, 418, 418, 418, 419, 419, 419, 419, 419, 419, 419, 419, 416, 416, 416, 416, 416, 416, 416, 416, 417, 417, 417, 417, 417, 417, 417, 417, 422, 422, 422, 422, 422, 422, 422, 422, 423, 423, 423, 423, 423, 423, 423, 423, 420, 420, 420, 420, 420, 420, 420, 420, 421, 421, 421, 421, 421, 421, 421, 421, -1
+// CHECK: 418, 418, 418, 418, 418, 418, 418, 418, 419, 419, 419, 419, 419, 419, 419, 419, 416, 416, 416, 416, 416, 416, 416, 416, 417, 417, 417, 417, 417, 417, 417, 417, 422, 422, 422, 422, 422, 422, 422, 422, 423, 423, 423, 423, 423, 423, 423, 423, 420, 420, 420, 420, 420, 420, 420, 420, 421, 421, 421, 421, 421, 421, 421, 421, 435, 435, 435, 435, 435, 435, 435, 435, 434, 434, 434, 434, 434, 434, 434, 434, 433, 433, 433, 433, 433, 433, 433, 433, 432, 432, 432, 432, 432, 432, 432, 432, 439, 439, 439, 439, 439, 439, 439, 439, 438, 438, 438, 438, 438, 438, 438, 438, 437, 437, 437, 437, 437, 437, 437, 437, 436, 436, 436, 436, 436, 436, 436, 436, -1
+// CHECK: 435, 435, 435, 435, 435, 435, 435, 435, 434, 434, 434, 434, 434, 434, 434, 434, 433, 433, 433, 433, 433, 433, 433, 433, 432, 432, 432, 432, 432, 432, 432, 432, 439, 439, 439, 439, 439, 439, 439, 439, 438, 438, 438, 438, 438, 438, 438, 438, 437, 437, 437, 437, 437, 437, 437, 437, 436, 436, 436, 436, 436, 436, 436, 436, 452, 452, 452, 452, 452, 452, 452, 452, 453, 453, 453, 453, 453, 453, 453, 453, 454, 454, 454, 454, 454, 454, 454, 454, 455, 455, 455, 455, 455, 455, 455, 455, 448, 448, 448, 448, 448, 448, 448, 448, 449, 449, 449, 449, 449, 449, 449, 449, 450, 450, 450, 450, 450, 450, 450, 450, 451, 451, 451, 451, 451, 451, 451, 451, -1
+// CHECK: 452, 452, 452, 452, 452, 452, 452, 452, 453, 453, 453, 453, 453, 453, 453, 453, 454, 454, 454, 454, 454, 454, 454, 454, 455, 455, 455, 455, 455, 455, 455, 455, 448, 448, 448, 448, 448, 448, 448, 448, 449, 449, 449, 449, 449, 449, 449, 449, 450, 450, 450, 450, 450, 450, 450, 450, 451, 451, 451, 451, 451, 451, 451, 451, 469, 469, 469, 469, 469, 469, 469, 469, 468, 468, 468, 468, 468, 468, 468, 468, 471, 471, 471, 471, 471, 471, 471, 471, 470, 470, 470, 470, 470, 470, 470, 470, 465, 465, 465, 465, 465, 465, 465, 465, 464, 464, 464, 464, 464, 464, 464, 464, 467, 467, 467, 467, 467, 467, 467, 467, 466, 466, 466, 466, 466, 466, 466, 466, -1
+// CHECK: 469, 469, 469, 469, 469, 469, 469, 469, 468, 468, 468, 468, 468, 468, 468, 468, 471, 471, 471, 471, 471, 471, 471, 471, 470, 470, 470, 470, 470, 470, 470, 470, 465, 465, 465, 465, 465, 465, 465, 465, 464, 464, 464, 464, 464, 464, 464, 464, 467, 467, 467, 467, 467, 467, 467, 467, 466, 466, 466, 466, 466, 466, 466, 466, 486, 486, 486, 486, 486, 486, 486, 486, 487, 487, 487, 487, 487, 487, 487, 487, 484, 484, 484, 484, 484, 484, 484, 484, 485, 485, 485, 485, 485, 485, 485, 485, 482, 482, 482, 482, 482, 482, 482, 482, 483, 483, 483, 483, 483, 483, 483, 483, 480, 480, 480, 480, 480, 480, 480, 480, 481, 481, 481, 481, 481, 481, 481, 481, -1
+// CHECK: 486, 486, 486, 486, 486, 486, 486, 486, 487, 487, 487, 487, 487, 487, 487, 487, 484, 484, 484, 484, 484, 484, 484, 484, 485, 485, 485, 485, 485, 485, 485, 485, 482, 482, 482, 482, 482, 482, 482, 482, 483, 483, 483, 483, 483, 483, 483, 483, 480, 480, 480, 480, 480, 480, 480, 480, 481, 481, 481, 481, 481, 481, 481, 481, 503, 503, 503, 503, 503, 503, 503, 503, 502, 502, 502, 502, 502, 502, 502, 502, 501, 501, 501, 501, 501, 501, 501, 501, 500, 500, 500, 500, 500, 500, 500, 500, 499, 499, 499, 499, 499, 499, 499, 499, 498, 498, 498, 498, 498, 498, 498, 498, 497, 497, 497, 497, 497, 497, 497, 497, 496, 496, 496, 496, 496, 496, 496, 496, -1
+// CHECK: 503, 503, 503, 503, 503, 503, 503, 503, 502, 502, 502, 502, 502, 502, 502, 502, 501, 501, 501, 501, 501, 501, 501, 501, 500, 500, 500, 500, 500, 500, 500, 500, 499, 499, 499, 499, 499, 499, 499, 499, 498, 498, 498, 498, 498, 498, 498, 498, 497, 497, 497, 497, 497, 497, 497, 497, 496, 496, 496, 496, 496, 496, 496, 496, 512, 512, 512, 512, 512, 512, 512, 512, 513, 513, 513, 513, 513, 513, 513, 513, 514, 514, 514, 514, 514, 514, 514, 514, 515, 515, 515, 515, 515, 515, 515, 515, 516, 516, 516, 516, 516, 516, 516, 516, 517, 517, 517, 517, 517, 517, 517, 517, 518, 518, 518, 518, 518, 518, 518, 518, 519, 519, 519, 519, 519, 519, 519, 519, -1
+// CHECK: 512, 512, 512, 512, 512, 512, 512, 512, 513, 513, 513, 513, 513, 513, 513, 513, 514, 514, 514, 514, 514, 514, 514, 514, 515, 515, 515, 515, 515, 515, 515, 515, 516, 516, 516, 516, 516, 516, 516, 516, 517, 517, 517, 517, 517, 517, 517, 517, 518, 518, 518, 518, 518, 518, 518, 518, 519, 519, 519, 519, 519, 519, 519, 519, 529, 529, 529, 529, 529, 529, 529, 529, 528, 528, 528, 528, 528, 528, 528, 528, 531, 531, 531, 531, 531, 531, 531, 531, 530, 530, 530, 530, 530, 530, 530, 530, 533, 533, 533, 533, 533, 533, 533, 533, 532, 532, 532, 532, 532, 532, 532, 532, 535, 535, 535, 535, 535, 535, 535, 535, 534, 534, 534, 534, 534, 534, 534, 534, -1
+// CHECK: 529, 529, 529, 529, 529, 529, 529, 529, 528, 528, 528, 528, 528, 528, 528, 528, 531, 531, 531, 531, 531, 531, 531, 531, 530, 530, 530, 530, 530, 530, 530, 530, 533, 533, 533, 533, 533, 533, 533, 533, 532, 532, 532, 532, 532, 532, 532, 532, 535, 535, 535, 535, 535, 535, 535, 535, 534, 534, 534, 534, 534, 534, 534, 534, 546, 546, 546, 546, 546, 546, 546, 546, 547, 547, 547, 547, 547, 547, 547, 547, 544, 544, 544, 544, 544, 544, 544, 544, 545, 545, 545, 545, 545, 545, 545, 545, 550, 550, 550, 550, 550, 550, 550, 550, 551, 551, 551, 551, 551, 551, 551, 551, 548, 548, 548, 548, 548, 548, 548, 548, 549, 549, 549, 549, 549, 549, 549, 549, -1
+// CHECK: 546, 546, 546, 546, 546, 546, 546, 546, 547, 547, 547, 547, 547, 547, 547, 547, 544, 544, 544, 544, 544, 544, 544, 544, 545, 545, 545, 545, 545, 545, 545, 545, 550, 550, 550, 550, 550, 550, 550, 550, 551, 551, 551, 551, 551, 551, 551, 551, 548, 548, 548, 548, 548, 548, 548, 548, 549, 549, 549, 549, 549, 549, 549, 549, 563, 563, 563, 563, 563, 563, 563, 563, 562, 562, 562, 562, 562, 562, 562, 562, 561, 561, 561, 561, 561, 561, 561, 561, 560, 560, 560, 560, 560, 560, 560, 560, 567, 567, 567, 567, 567, 567, 567, 567, 566, 566, 566, 566, 566, 566, 566, 566, 565, 565, 565, 565, 565, 565, 565, 565, 564, 564, 564, 564, 564, 564, 564, 564, -1
+// CHECK: 563, 563, 563, 563, 563, 563, 563, 563, 562, 562, 562, 562, 562, 562, 562, 562, 561, 561, 561, 561, 561, 561, 561, 561, 560, 560, 560, 560, 560, 560, 560, 560, 567, 567, 567, 567, 567, 567, 567, 567, 566, 566, 566, 566, 566, 566, 566, 566, 565, 565, 565, 565, 565, 565, 565, 565, 564, 564, 564, 564, 564, 564, 564, 564, 580, 580, 580, 580, 580, 580, 580, 580, 581, 581, 581, 581, 581, 581, 581, 581, 582, 582, 582, 582, 582, 582, 582, 582, 583, 583, 583, 583, 583, 583, 583, 583, 576, 576, 576, 576, 576, 576, 576, 576, 577, 577, 577, 577, 577, 577, 577, 577, 578, 578, 578, 578, 578, 578, 578, 578, 579, 579, 579, 579, 579, 579, 579, 579, -1
+// CHECK: 580, 580, 580, 580, 580, 580, 580, 580, 581, 581, 581, 581, 581, 581, 581, 581, 582, 582, 582, 582, 582, 582, 582, 582, 583, 583, 583, 583, 583, 583, 583, 583, 576, 576, 576, 576, 576, 576, 576, 576, 577, 577, 577, 577, 577, 577, 577, 577, 578, 578, 578, 578, 578, 578, 578, 578, 579, 579, 579, 579, 579, 579, 579, 579, 597, 597, 597, 597, 597, 597, 597, 597, 596, 596, 596, 596, 596, 596, 596, 596, 599, 599, 599, 599, 599, 599, 599, 599, 598, 598, 598, 598, 598, 598, 598, 598, 593, 593, 593, 593, 593, 593, 593, 593, 592, 592, 592, 592, 592, 592, 592, 592, 595, 595, 595, 595, 595, 595, 595, 595, 594, 594, 594, 594, 594, 594, 594, 594, -1
+// CHECK: 597, 597, 597, 597, 597, 597, 597, 597, 596, 596, 596, 596, 596, 596, 596, 596, 599, 599, 599, 599, 599, 599, 599, 599, 598, 598, 598, 598, 598, 598, 598, 598, 593, 593, 593, 593, 593, 593, 593, 593, 592, 592, 592, 592, 592, 592, 592, 592, 595, 595, 595, 595, 595, 595, 595, 595, 594, 594, 594, 594, 594, 594, 594, 594, 614, 614, 614, 614, 614, 614, 614, 614, 615, 615, 615, 615, 615, 615, 615, 615, 612, 612, 612, 612, 612, 612, 612, 612, 613, 613, 613, 613, 613, 613, 613, 613, 610, 610, 610, 610, 610, 610, 610, 610, 611, 611, 611, 611, 611, 611, 611, 611, 608, 608, 608, 608, 608, 608, 608, 608, 609, 609, 609, 609, 609, 609, 609, 609, -1
+// CHECK: 614, 614, 614, 614, 614, 614, 614, 614, 615, 615, 615, 615, 615, 615, 615, 615, 612, 612, 612, 612, 612, 612, 612, 612, 613, 613, 613, 613, 613, 613, 613, 613, 610, 610, 610, 610, 610, 610, 610, 610, 611, 611, 611, 611, 611, 611, 611, 611, 608, 608, 608, 608, 608, 608, 608, 608, 609, 609, 609, 609, 609, 609, 609, 609, 631, 631, 631, 631, 631, 631, 631, 631, 630, 630, 630, 630, 630, 630, 630, 630, 629, 629, 629, 629, 629, 629, 629, 629, 628, 628, 628, 628, 628, 628, 628, 628, 627, 627, 627, 627, 627, 627, 627, 627, 626, 626, 626, 626, 626, 626, 626, 626, 625, 625, 625, 625, 625, 625, 625, 625, 624, 624, 624, 624, 624, 624, 624, 624, -1
+// CHECK: 631, 631, 631, 631, 631, 631, 631, 631, 630, 630, 630, 630, 630, 630, 630, 630, 629, 629, 629, 629, 629, 629, 629, 629, 628, 628, 628, 628, 628, 628, 628, 628, 627, 627, 627, 627, 627, 627, 627, 627, 626, 626, 626, 626, 626, 626, 626, 626, 625, 625, 625, 625, 625, 625, 625, 625, 624, 624, 624, 624, 624, 624, 624, 624, 640, 640, 640, 640, 640, 640, 640, 640, 641, 641, 641, 641, 641, 641, 641, 641, 642, 642, 642, 642, 642, 642, 642, 642, 643, 643, 643, 643, 643, 643, 643, 643, 644, 644, 644, 644, 644, 644, 644, 644, 645, 645, 645, 645, 645, 645, 645, 645, 646, 646, 646, 646, 646, 646, 646, 646, 647, 647, 647, 647, 647, 647, 647, 647, -1
+// CHECK: 640, 640, 640, 640, 640, 640, 640, 640, 641, 641, 641, 641, 641, 641, 641, 641, 642, 642, 642, 642, 642, 642, 642, 642, 643, 643, 643, 643, 643, 643, 643, 643, 644, 644, 644, 644, 644, 644, 644, 644, 645, 645, 645, 645, 645, 645, 645, 645, 646, 646, 646, 646, 646, 646, 646, 646, 647, 647, 647, 647, 647, 647, 647, 647, 657, 657, 657, 657, 657, 657, 657, 657, 656, 656, 656, 656, 656, 656, 656, 656, 659, 659, 659, 659, 659, 659, 659, 659, 658, 658, 658, 658, 658, 658, 658, 658, 661, 661, 661, 661, 661, 661, 661, 661, 660, 660, 660, 660, 660, 660, 660, 660, 663, 663, 663, 663, 663, 663, 663, 663, 662, 662, 662, 662, 662, 662, 662, 662, -1
+// CHECK: 657, 657, 657, 657, 657, 657, 657, 657, 656, 656, 656, 656, 656, 656, 656, 656, 659, 659, 659, 659, 659, 659, 659, 659, 658, 658, 658, 658, 658, 658, 658, 658, 661, 661, 661, 661, 661, 661, 661, 661, 660, 660, 660, 660, 660, 660, 660, 660, 663, 663, 663, 663, 663, 663, 663, 663, 662, 662, 662, 662, 662, 662, 662, 662, 674, 674, 674, 674, 674, 674, 674, 674, 675, 675, 675, 675, 675, 675, 675, 675, 672, 672, 672, 672, 672, 672, 672, 672, 673, 673, 673, 673, 673, 673, 673, 673, 678, 678, 678, 678, 678, 678, 678, 678, 679, 679, 679, 679, 679, 679, 679, 679, 676, 676, 676, 676, 676, 676, 676, 676, 677, 677, 677, 677, 677, 677, 677, 677, -1
+// CHECK: 674, 674, 674, 674, 674, 674, 674, 674, 675, 675, 675, 675, 675, 675, 675, 675, 672, 672, 672, 672, 672, 672, 672, 672, 673, 673, 673, 673, 673, 673, 673, 673, 678, 678, 678, 678, 678, 678, 678, 678, 679, 679, 679, 679, 679, 679, 679, 679, 676, 676, 676, 676, 676, 676, 676, 676, 677, 677, 677, 677, 677, 677, 677, 677, 691, 691, 691, 691, 691, 691, 691, 691, 690, 690, 690, 690, 690, 690, 690, 690, 689, 689, 689, 689, 689, 689, 689, 689, 688, 688, 688, 688, 688, 688, 688, 688, 695, 695, 695, 695, 695, 695, 695, 695, 694, 694, 694, 694, 694, 694, 694, 694, 693, 693, 693, 693, 693, 693, 693, 693, 692, 692, 692, 692, 692, 692, 692, 692, -1
+// CHECK: 691, 691, 691, 691, 691, 691, 691, 691, 690, 690, 690, 690, 690, 690, 690, 690, 689, 689, 689, 689, 689, 689, 689, 689, 688, 688, 688, 688, 688, 688, 688, 688, 695, 695, 695, 695, 695, 695, 695, 695, 694, 694, 694, 694, 694, 694, 694, 694, 693, 693, 693, 693, 693, 693, 693, 693, 692, 692, 692, 692, 692, 692, 692, 692, 708, 708, 708, 708, 708, 708, 708, 708, 709, 709, 709, 709, 709, 709, 709, 709, 710, 710, 710, 710, 710, 710, 710, 710, 711, 711, 711, 711, 711, 711, 711, 711, 704, 704, 704, 704, 704, 704, 704, 704, 705, 705, 705, 705, 705, 705, 705, 705, 706, 706, 706, 706, 706, 706, 706, 706, 707, 707, 707, 707, 707, 707, 707, 707, -1
+// CHECK: 708, 708, 708, 708, 708, 708, 708, 708, 709, 709, 709, 709, 709, 709, 709, 709, 710, 710, 710, 710, 710, 710, 710, 710, 711, 711, 711, 711, 711, 711, 711, 711, 704, 704, 704, 704, 704, 704, 704, 704, 705, 705, 705, 705, 705, 705, 705, 705, 706, 706, 706, 706, 706, 706, 706, 706, 707, 707, 707, 707, 707, 707, 707, 707, 725, 725, 725, 725, 725, 725, 725, 725, 724, 724, 724, 724, 724, 724, 724, 724, 727, 727, 727, 727, 727, 727, 727, 727, 726, 726, 726, 726, 726, 726, 726, 726, 721, 721, 721, 721, 721, 721, 721, 721, 720, 720, 720, 720, 720, 720, 720, 720, 723, 723, 723, 723, 723, 723, 723, 723, 722, 722, 722, 722, 722, 722, 722, 722, -1
+// CHECK: 725, 725, 725, 725, 725, 725, 725, 725, 724, 724, 724, 724, 724, 724, 724, 724, 727, 727, 727, 727, 727, 727, 727, 727, 726, 726, 726, 726, 726, 726, 726, 726, 721, 721, 721, 721, 721, 721, 721, 721, 720, 720, 720, 720, 720, 720, 720, 720, 723, 723, 723, 723, 723, 723, 723, 723, 722, 722, 722, 722, 722, 722, 722, 722, 742, 742, 742, 742, 742, 742, 742, 742, 743, 743, 743, 743, 743, 743, 743, 743, 740, 740, 740, 740, 740, 740, 740, 740, 741, 741, 741, 741, 741, 741, 741, 741, 738, 738, 738, 738, 738, 738, 738, 738, 739, 739, 739, 739, 739, 739, 739, 739, 736, 736, 736, 736, 736, 736, 736, 736, 737, 737, 737, 737, 737, 737, 737, 737, -1
+// CHECK: 742, 742, 742, 742, 742, 742, 742, 742, 743, 743, 743, 743, 743, 743, 743, 743, 740, 740, 740, 740, 740, 740, 740, 740, 741, 741, 741, 741, 741, 741, 741, 741, 738, 738, 738, 738, 738, 738, 738, 738, 739, 739, 739, 739, 739, 739, 739, 739, 736, 736, 736, 736, 736, 736, 736, 736, 737, 737, 737, 737, 737, 737, 737, 737, 759, 759, 759, 759, 759, 759, 759, 759, 758, 758, 758, 758, 758, 758, 758, 758, 757, 757, 757, 757, 757, 757, 757, 757, 756, 756, 756, 756, 756, 756, 756, 756, 755, 755, 755, 755, 755, 755, 755, 755, 754, 754, 754, 754, 754, 754, 754, 754, 753, 753, 753, 753, 753, 753, 753, 753, 752, 752, 752, 752, 752, 752, 752, 752, -1
+// CHECK: 759, 759, 759, 759, 759, 759, 759, 759, 758, 758, 758, 758, 758, 758, 758, 758, 757, 757, 757, 757, 757, 757, 757, 757, 756, 756, 756, 756, 756, 756, 756, 756, 755, 755, 755, 755, 755, 755, 755, 755, 754, 754, 754, 754, 754, 754, 754, 754, 753, 753, 753, 753, 753, 753, 753, 753, 752, 752, 752, 752, 752, 752, 752, 752, 768, 768, 768, 768, 768, 768, 768, 768, 769, 769, 769, 769, 769, 769, 769, 769, 770, 770, 770, 770, 770, 770, 770, 770, 771, 771, 771, 771, 771, 771, 771, 771, 772, 772, 772, 772, 772, 772, 772, 772, 773, 773, 773, 773, 773, 773, 773, 773, 774, 774, 774, 774, 774, 774, 774, 774, 775, 775, 775, 775, 775, 775, 775, 775, -1
+// CHECK: 768, 768, 768, 768, 768, 768, 768, 768, 769, 769, 769, 769, 769, 769, 769, 769, 770, 770, 770, 770, 770, 770, 770, 770, 771, 771, 771, 771, 771, 771, 771, 771, 772, 772, 772, 772, 772, 772, 772, 772, 773, 773, 773, 773, 773, 773, 773, 773, 774, 774, 774, 774, 774, 774, 774, 774, 775, 775, 775, 775, 775, 775, 775, 775, 785, 785, 785, 785, 785, 785, 785, 785, 784, 784, 784, 784, 784, 784, 784, 784, 787, 787, 787, 787, 787, 787, 787, 787, 786, 786, 786, 786, 786, 786, 786, 786, 789, 789, 789, 789, 789, 789, 789, 789, 788, 788, 788, 788, 788, 788, 788, 788, 791, 791, 791, 791, 791, 791, 791, 791, 790, 790, 790, 790, 790, 790, 790, 790, -1
+// CHECK: 785, 785, 785, 785, 785, 785, 785, 785, 784, 784, 784, 784, 784, 784, 784, 784, 787, 787, 787, 787, 787, 787, 787, 787, 786, 786, 786, 786, 786, 786, 786, 786, 789, 789, 789, 789, 789, 789, 789, 789, 788, 788, 788, 788, 788, 788, 788, 788, 791, 791, 791, 791, 791, 791, 791, 791, 790, 790, 790, 790, 790, 790, 790, 790, 802, 802, 802, 802, 802, 802, 802, 802, 803, 803, 803, 803, 803, 803, 803, 803, 800, 800, 800, 800, 800, 800, 800, 800, 801, 801, 801, 801, 801, 801, 801, 801, 806, 806, 806, 806, 806, 806, 806, 806, 807, 807, 807, 807, 807, 807, 807, 807, 804, 804, 804, 804, 804, 804, 804, 804, 805, 805, 805, 805, 805, 805, 805, 805, -1
+// CHECK: 802, 802, 802, 802, 802, 802, 802, 802, 803, 803, 803, 803, 803, 803, 803, 803, 800, 800, 800, 800, 800, 800, 800, 800, 801, 801, 801, 801, 801, 801, 801, 801, 806, 806, 806, 806, 806, 806, 806, 806, 807, 807, 807, 807, 807, 807, 807, 807, 804, 804, 804, 804, 804, 804, 804, 804, 805, 805, 805, 805, 805, 805, 805, 805, 819, 819, 819, 819, 819, 819, 819, 819, 818, 818, 818, 818, 818, 818, 818, 818, 817, 817, 817, 817, 817, 817, 817, 817, 816, 816, 816, 816, 816, 816, 816, 816, 823, 823, 823, 823, 823, 823, 823, 823, 822, 822, 822, 822, 822, 822, 822, 822, 821, 821, 821, 821, 821, 821, 821, 821, 820, 820, 820, 820, 820, 820, 820, 820, -1
+// CHECK: 819, 819, 819, 819, 819, 819, 819, 819, 818, 818, 818, 818, 818, 818, 818, 818, 817, 817, 817, 817, 817, 817, 817, 817, 816, 816, 816, 816, 816, 816, 816, 816, 823, 823, 823, 823, 823, 823, 823, 823, 822, 822, 822, 822, 822, 822, 822, 822, 821, 821, 821, 821, 821, 821, 821, 821, 820, 820, 820, 820, 820, 820, 820, 820, 836, 836, 836, 836, 836, 836, 836, 836, 837, 837, 837, 837, 837, 837, 837, 837, 838, 838, 838, 838, 838, 838, 838, 838, 839, 839, 839, 839, 839, 839, 839, 839, 832, 832, 832, 832, 832, 832, 832, 832, 833, 833, 833, 833, 833, 833, 833, 833, 834, 834, 834, 834, 834, 834, 834, 834, 835, 835, 835, 835, 835, 835, 835, 835, -1
+// CHECK: 836, 836, 836, 836, 836, 836, 836, 836, 837, 837, 837, 837, 837, 837, 837, 837, 838, 838, 838, 838, 838, 838, 838, 838, 839, 839, 839, 839, 839, 839, 839, 839, 832, 832, 832, 832, 832, 832, 832, 832, 833, 833, 833, 833, 833, 833, 833, 833, 834, 834, 834, 834, 834, 834, 834, 834, 835, 835, 835, 835, 835, 835, 835, 835, 853, 853, 853, 853, 853, 853, 853, 853, 852, 852, 852, 852, 852, 852, 852, 852, 855, 855, 855, 855, 855, 855, 855, 855, 854, 854, 854, 854, 854, 854, 854, 854, 849, 849, 849, 849, 849, 849, 849, 849, 848, 848, 848, 848, 848, 848, 848, 848, 851, 851, 851, 851, 851, 851, 851, 851, 850, 850, 850, 850, 850, 850, 850, 850, -1
+// CHECK: 853, 853, 853, 853, 853, 853, 853, 853, 852, 852, 852, 852, 852, 852, 852, 852, 855, 855, 855, 855, 855, 855, 855, 855, 854, 854, 854, 854, 854, 854, 854, 854, 849, 849, 849, 849, 849, 849, 849, 849, 848, 848, 848, 848, 848, 848, 848, 848, 851, 851, 851, 851, 851, 851, 851, 851, 850, 850, 850, 850, 850, 850, 850, 850, 870, 870, 870, 870, 870, 870, 870, 870, 871, 871, 871, 871, 871, 871, 871, 871, 868, 868, 868, 868, 868, 868, 868, 868, 869, 869, 869, 869, 869, 869, 869, 869, 866, 866, 866, 866, 866, 866, 866, 866, 867, 867, 867, 867, 867, 867, 867, 867, 864, 864, 864, 864, 864, 864, 864, 864, 865, 865, 865, 865, 865, 865, 865, 865, -1
+// CHECK: 870, 870, 870, 870, 870, 870, 870, 870, 871, 871, 871, 871, 871, 871, 871, 871, 868, 868, 868, 868, 868, 868, 868, 868, 869, 869, 869, 869, 869, 869, 869, 869, 866, 866, 866, 866, 866, 866, 866, 866, 867, 867, 867, 867, 867, 867, 867, 867, 864, 864, 864, 864, 864, 864, 864, 864, 865, 865, 865, 865, 865, 865, 865, 865, 887, 887, 887, 887, 887, 887, 887, 887, 886, 886, 886, 886, 886, 886, 886, 886, 885, 885, 885, 885, 885, 885, 885, 885, 884, 884, 884, 884, 884, 884, 884, 884, 883, 883, 883, 883, 883, 883, 883, 883, 882, 882, 882, 882, 882, 882, 882, 882, 881, 881, 881, 881, 881, 881, 881, 881, 880, 880, 880, 880, 880, 880, 880, 880, -1
+// CHECK: 887, 887, 887, 887, 887, 887, 887, 887, 886, 886, 886, 886, 886, 886, 886, 886, 885, 885, 885, 885, 885, 885, 885, 885, 884, 884, 884, 884, 884, 884, 884, 884, 883, 883, 883, 883, 883, 883, 883, 883, 882, 882, 882, 882, 882, 882, 882, 882, 881, 881, 881, 881, 881, 881, 881, 881, 880, 880, 880, 880, 880, 880, 880, 880, 896, 896, 896, 896, 896, 896, 896, 896, 897, 897, 897, 897, 897, 897, 897, 897, 898, 898, 898, 898, 898, 898, 898, 898, 899, 899, 899, 899, 899, 899, 899, 899, 900, 900, 900, 900, 900, 900, 900, 900, 901, 901, 901, 901, 901, 901, 901, 901, 902, 902, 902, 902, 902, 902, 902, 902, 903, 903, 903, 903, 903, 903, 903, 903, -1
+// CHECK: 896, 896, 896, 896, 896, 896, 896, 896, 897, 897, 897, 897, 897, 897, 897, 897, 898, 898, 898, 898, 898, 898, 898, 898, 899, 899, 899, 899, 899, 899, 899, 899, 900, 900, 900, 900, 900, 900, 900, 900, 901, 901, 901, 901, 901, 901, 901, 901, 902, 902, 902, 902, 902, 902, 902, 902, 903, 903, 903, 903, 903, 903, 903, 903, 913, 913, 913, 913, 913, 913, 913, 913, 912, 912, 912, 912, 912, 912, 912, 912, 915, 915, 915, 915, 915, 915, 915, 915, 914, 914, 914, 914, 914, 914, 914, 914, 917, 917, 917, 917, 917, 917, 917, 917, 916, 916, 916, 916, 916, 916, 916, 916, 919, 919, 919, 919, 919, 919, 919, 919, 918, 918, 918, 918, 918, 918, 918, 918, -1
+// CHECK: 913, 913, 913, 913, 913, 913, 913, 913, 912, 912, 912, 912, 912, 912, 912, 912, 915, 915, 915, 915, 915, 915, 915, 915, 914, 914, 914, 914, 914, 914, 914, 914, 917, 917, 917, 917, 917, 917, 917, 917, 916, 916, 916, 916, 916, 916, 916, 916, 919, 919, 919, 919, 919, 919, 919, 919, 918, 918, 918, 918, 918, 918, 918, 918, 930, 930, 930, 930, 930, 930, 930, 930, 931, 931, 931, 931, 931, 931, 931, 931, 928, 928, 928, 928, 928, 928, 928, 928, 929, 929, 929, 929, 929, 929, 929, 929, 934, 934, 934, 934, 934, 934, 934, 934, 935, 935, 935, 935, 935, 935, 935, 935, 932, 932, 932, 932, 932, 932, 932, 932, 933, 933, 933, 933, 933, 933, 933, 933, -1
+// CHECK: 930, 930, 930, 930, 930, 930, 930, 930, 931, 931, 931, 931, 931, 931, 931, 931, 928, 928, 928, 928, 928, 928, 928, 928, 929, 929, 929, 929, 929, 929, 929, 929, 934, 934, 934, 934, 934, 934, 934, 934, 935, 935, 935, 935, 935, 935, 935, 935, 932, 932, 932, 932, 932, 932, 932, 932, 933, 933, 933, 933, 933, 933, 933, 933, 947, 947, 947, 947, 947, 947, 947, 947, 946, 946, 946, 946, 946, 946, 946, 946, 945, 945, 945, 945, 945, 945, 945, 945, 944, 944, 944, 944, 944, 944, 944, 944, 951, 951, 951, 951, 951, 951, 951, 951, 950, 950, 950, 950, 950, 950, 950, 950, 949, 949, 949, 949, 949, 949, 949, 949, 948, 948, 948, 948, 948, 948, 948, 948, -1
+// CHECK: 947, 947, 947, 947, 947, 947, 947, 947, 946, 946, 946, 946, 946, 946, 946, 946, 945, 945, 945, 945, 945, 945, 945, 945, 944, 944, 944, 944, 944, 944, 944, 944, 951, 951, 951, 951, 951, 951, 951, 951, 950, 950, 950, 950, 950, 950, 950, 950, 949, 949, 949, 949, 949, 949, 949, 949, 948, 948, 948, 948, 948, 948, 948, 948, 964, 964, 964, 964, 964, 964, 964, 964, 965, 965, 965, 965, 965, 965, 965, 965, 966, 966, 966, 966, 966, 966, 966, 966, 967, 967, 967, 967, 967, 967, 967, 967, 960, 960, 960, 960, 960, 960, 960, 960, 961, 961, 961, 961, 961, 961, 961, 961, 962, 962, 962, 962, 962, 962, 962, 962, 963, 963, 963, 963, 963, 963, 963, 963, -1
+// CHECK: 964, 964, 964, 964, 964, 964, 964, 964, 965, 965, 965, 965, 965, 965, 965, 965, 966, 966, 966, 966, 966, 966, 966, 966, 967, 967, 967, 967, 967, 967, 967, 967, 960, 960, 960, 960, 960, 960, 960, 960, 961, 961, 961, 961, 961, 961, 961, 961, 962, 962, 962, 962, 962, 962, 962, 962, 963, 963, 963, 963, 963, 963, 963, 963, 981, 981, 981, 981, 981, 981, 981, 981, 980, 980, 980, 980, 980, 980, 980, 980, 983, 983, 983, 983, 983, 983, 983, 983, 982, 982, 982, 982, 982, 982, 982, 982, 977, 977, 977, 977, 977, 977, 977, 977, 976, 976, 976, 976, 976, 976, 976, 976, 979, 979, 979, 979, 979, 979, 979, 979, 978, 978, 978, 978, 978, 978, 978, 978, -1
+// CHECK: 981, 981, 981, 981, 981, 981, 981, 981, 980, 980, 980, 980, 980, 980, 980, 980, 983, 983, 983, 983, 983, 983, 983, 983, 982, 982, 982, 982, 982, 982, 982, 982, 977, 977, 977, 977, 977, 977, 977, 977, 976, 976, 976, 976, 976, 976, 976, 976, 979, 979, 979, 979, 979, 979, 979, 979, 978, 978, 978, 978, 978, 978, 978, 978, 998, 998, 998, 998, 998, 998, 998, 998, 999, 999, 999, 999, 999, 999, 999, 999, 996, 996, 996, 996, 996, 996, 996, 996, 997, 997, 997, 997, 997, 997, 997, 997, 994, 994, 994, 994, 994, 994, 994, 994, 995, 995, 995, 995, 995, 995, 995, 995, 992, 992, 992, 992, 992, 992, 992, 992, 993, 993, 993, 993, 993, 993, 993, 993, -1
+// CHECK: 998, 998, 998, 998, 998, 998, 998, 998, 999, 999, 999, 999, 999, 999, 999, 999, 996, 996, 996, 996, 996, 996, 996, 996, 997, 997, 997, 997, 997, 997, 997, 997, 994, 994, 994, 994, 994, 994, 994, 994, 995, 995, 995, 995, 995, 995, 995, 995, 992, 992, 992, 992, 992, 992, 992, 992, 993, 993, 993, 993, 993, 993, 993, 993, 1015, 1015, 1015, 1015, 1015, 1015, 1015, 1015, 1014, 1014, 1014, 1014, 1014, 1014, 1014, 1014, 1013, 1013, 1013, 1013, 1013, 1013, 1013, 1013, 1012, 1012, 1012, 1012, 1012, 1012, 1012, 1012, 1011, 1011, 1011, 1011, 1011, 1011, 1011, 1011, 1010, 1010, 1010, 1010, 1010, 1010, 1010, 1010, 1009, 1009, 1009, 1009, 1009, 1009, 1009, 1009, 1008, 1008, 1008, 1008, 1008, 1008, 1008, 1008, -1
+// CHECK: 1015, 1015, 1015, 1015, 1015, 1015, 1015, 1015, 1014, 1014, 1014, 1014, 1014, 1014, 1014, 1014, 1013, 1013, 1013, 1013, 1013, 1013, 1013, 1013, 1012, 1012, 1012, 1012, 1012, 1012, 1012, 1012, 1011, 1011, 1011, 1011, 1011, 1011, 1011, 1011, 1010, 1010, 1010, 1010, 1010, 1010, 1010, 1010, 1009, 1009, 1009, 1009, 1009, 1009, 1009, 1009, 1008, 1008, 1008, 1008, 1008, 1008, 1008, 1008, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -1
+
>From 2689497420af79e02b9fd54287c13ffc54a4fe57 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.ozen at gmail.com>
Date: Fri, 15 Sep 2023 09:29:54 +0200
Subject: [PATCH 2/2] use mlir-cpu-runner and make the name of the test
consistent
---
...4_swizzle128b.mlir => tma_load_64x64_swizzle128b.mlir} | 8 +++++---
1 file changed, 5 insertions(+), 3 deletions(-)
rename mlir/test/Integration/GPU/CUDA/sm90/{tmaload_64_64_swizzle128b.mlir => tma_load_64x64_swizzle128b.mlir} (99%)
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
similarity index 99%
rename from mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir
rename to mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
index b3c55c87405c507..140b0c2ecd77aa7 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/tmaload_64_64_swizzle128b.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
@@ -27,9 +27,11 @@
// 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
+// RUN: | mlir-cpu-runner \
+// RUN: --shared-libs=%mlir_cuda_runtime \
+// RUN: --shared-libs=%mlir_runner_utils \
+// RUN: --entry-point-result=void \
+// RUN: | FileCheck %s
// This does 3 TMA loads with 128B Swizzling :
// TMA Load: Matrix-A[0:128][0:64]
More information about the Mlir-commits
mailing list