Skip to content

[mlir] Fixed typo in type (128x64 -> 64x128) in TMA load test #70022

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Oct 26, 2023

Conversation

grypp
Copy link
Member

@grypp grypp commented Oct 24, 2023

The test was meant to check 64x128xf16 as the contiguous dimension exceeds the cache line (128b). TMA requires cache line-aligned loads, so loading 64x128 can be done with two 64x64 loads, as documented in the test.

However, there was a typo in the type, which was memref<128x64xf16> instead of the correct memref<64x128xf16>. This PR corrects the issue and updates the verification.

The test was meant to check `64x128xf16` as the contiguous dimension exceeds the cache line (128b). TMA requires cache line-aligned loads, so loading 64x128 can be done with two 64x64 loads, as documented in the test.

However, there was a typo in the type, which was `memref<128x64xf16>` instead of the correct `memref<64x128xf16>`. This PR corrects the issue and updates the verification.
@llvmbot
Copy link
Member

llvmbot commented Oct 24, 2023

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-mlir-gpu

Author: Guray Ozen (grypp)

Changes

The test was meant to check 64x128xf16 as the contiguous dimension exceeds the cache line (128b). TMA requires cache line-aligned loads, so loading 64x128 can be done with two 64x64 loads, as documented in the test.

However, there was a typo in the type, which was memref&lt;128x64xf16&gt; instead of the correct memref&lt;64x128xf16&gt;. This PR corrects the issue and updates the verification.


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

1 Files Affected:

  • (modified) mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir (+69-74)
diff --git a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
index 242c5ff875cf44a..13b9c48dabe85d7 100644
--- a/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
+++ b/mlir/test/Integration/GPU/CUDA/sm90/tma_load_64x64_swizzle128b.mlir
@@ -33,8 +33,8 @@
 !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>
+!rhs = memref<64x128xf16>
+!shmemrhs = memref<64x128xf16, 3>
 !rhsTensorMap = !nvgpu.tensormap.descriptor<tensor = !shmemrhs, swizzle = swizzle_128b, l2promo=none, oob=zero, interleave=none>
 
 module @mymod {
@@ -68,10 +68,6 @@ module @mymod {
         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
@@ -103,20 +99,19 @@ module @mymod {
       %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>
+      %rhsShmem2 = memref.subview %rhsShmem[32, 0][128, 64][1, 1] : !shmemrhs to memref<128x64xf16, strided<[128, 1], offset: 4096>, 3>
     
       // Step 5. Initialize the mbarrier
       %9 = nvgpu.mbarrier.create -> !barrierType
       nvgpu.mbarrier.init %9[%c0], %5 : !barrierType
       %10 = arith.cmpi eq, %6, %c0 : index
       
-      
       // Step 6. 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[%c0] to %lhsShmem : !lhsTensorMap, !barrierType -> !shmemlhs
         nvgpu.tma.async.load %d_rhsTensorMap[%c0, %c0], %9[%c0] to %rhsShmem : !rhsTensorMap, !barrierType -> !shmemrhs
-        nvgpu.tma.async.load %d_rhsTensorMap[%c64, %c0], %9[%c0] to %rhsShmem2 : !rhsTensorMap, !barrierType -> memref<?x?xf16, strided<[?, ?], offset: ?>, 3>
+        nvgpu.tma.async.load %d_rhsTensorMap[%c64, %c0], %9[%c0] to %rhsShmem2 : !rhsTensorMap, !barrierType -> memref<128x64xf16, strided<[128, 1], offset: 4096>, 3>
         nvgpu.mbarrier.arrive.expect_tx %9[%c0], %c32768 : !barrierType
       } else {
         nvgpu.mbarrier.arrive.expect_tx %9[%c0], %c0 : !barrierType
@@ -147,69 +142,69 @@ module @mymod {
 
 
 // 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...
[truncated]

Copy link
Collaborator

@qcolombet qcolombet left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Github chokes on the diff but based on the comment (fixing the dimensions order), sounds good.

@grypp grypp merged commit f7dc26c into llvm:main Oct 26, 2023
@grypp grypp deleted the fix-tma-test branch October 26, 2023 08:02
zahiraam pushed a commit to zahiraam/llvm-project that referenced this pull request Oct 26, 2023
…0022)

The test was meant to check `64x128xf16` as the contiguous dimension
exceeds the cache line (128b). TMA requires cache line-aligned loads, so
loading 64x128 can be done with two 64x64 loads, as documented in the
test.

However, there was a typo in the type, which was `memref<128x64xf16>`
instead of the correct `memref<64x128xf16>`. This PR corrects the issue
and updates the verification.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants