T.cuda.sm90.Wgmma
  target             nvidia.h200_sxm
  target capability  wgmma.mma_async
  execution mesh     Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),))))
                       predicates:
                         each top-level mode has no backward step
                         each top-level mode reaches each of its own slots exactly once
                         p0 % 128 == 0
  parameters
    n        any value
               predicates:
                 8 <= n <= 256
                 n % 8 == 0
    dtype    any value
               predicates:
                 dtype in {bf16, f16, fp8e4m3}
    form     any value
               predicates:
                 form in {Form.SS, Form.RS}
    a_major  dtype=bf16, form=SS  a_major in {Major.MN, Major.K} (default Major.MN)
             dtype=bf16, form=RS  Major.K
             dtype=f16, form=SS   a_major in {Major.MN, Major.K}
             dtype=f16, form=RS   Major.K
             dtype=fp8e4m3        Major.K
    b_major  dtype=bf16     b_major in {Major.MN, Major.K} (default Major.MN)
             dtype=f16      b_major in {Major.MN, Major.K}
             dtype=fp8e4m3  Major.K
  operands
    C  shape=(64, n) dtype=f32 storage=rmem, held in 1 arrangement:
         ShardLayout(Layout((4, 2, 8, c, 4, 2), (16, 8, 1, 512, 128, 64)), (S(0), S(2), S(4)), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
         predicates:
           each top-level mode has no backward step
           each top-level mode reaches each of its own slots exactly once
           p0 % 128 == 0
           8 * c == n
    A  dtype=bf16, form=SS     shape=(64, 16) dtype=bf16 storage=smem, held in 2 arrangements:
                                 a_major=MN  ShardLayout(ComposedLayout(Swizzle(a_swizzle, 4, 3), 0, Layout(((W, r), (8, 2)), ((a_stride_00, a_swizzle=0  s
                                             a_swizzle=1  l
                                             a_swizzle=2  l
                                             a_swizzle=3  l), (a_stride_10, a_swizzle=0  l
                                             a_swizzle=1  s
                                             a_swizzle=2  s
                                             a_swizzle=3  s)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                               predicates:
                                                 W == (8, 16, 32, 64)[a_swizzle]
                                                 r * W == 64
                                                 a_stride_00 == 1
                                                 a_stride_10 == W
                                                 l % 8 == 0
                                 a_major=K   ShardLayout(ComposedLayout(Swizzle(a_swizzle, 4, 3), k0, Layout(((8, r), (8, 2)), ((a_stride_00, s), (a_stride_10, a_stride_11)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                               predicates:
                                                 8 * r == 64
                                                 a_stride_00 == (8, 16, 32, 64)[a_swizzle]
                                                 a_stride_10 == 1
                                                 (a_swizzle == 0) | (a_stride_11 == 8)
                                                 a_stride_11 % 8 == 0
                                                 k0 % 16 == 0
                                                 k0 >= 0
                                                 k0 < (8, 16, 32, 64)[a_swizzle]
                                 predicates:
                                   s % 8 == 0
       dtype=bf16, form=RS     shape=(64, 16) dtype=bf16 storage=rmem, held in 1 arrangement:
                                 ShardLayout(Layout((4, 2, 8, 2, 4, 2), (16, 8, 1, 512, 128, 64)), (S(0), S(2), S(4)), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
       dtype=f16, form=SS      shape=(64, 16) dtype=f16 storage=smem, held in 2 arrangements:
                                 a_major=MN  ShardLayout(ComposedLayout(Swizzle(a_swizzle, 4, 3), 0, Layout(((W, r), (8, 2)), ((a_stride_00, a_swizzle=0  s
                                             a_swizzle=1  l
                                             a_swizzle=2  l
                                             a_swizzle=3  l), (a_stride_10, a_swizzle=0  l
                                             a_swizzle=1  s
                                             a_swizzle=2  s
                                             a_swizzle=3  s)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                               predicates:
                                                 W == (8, 16, 32, 64)[a_swizzle]
                                                 r * W == 64
                                                 a_stride_00 == 1
                                                 a_stride_10 == W
                                                 l % 8 == 0
                                 a_major=K   ShardLayout(ComposedLayout(Swizzle(a_swizzle, 4, 3), k0, Layout(((8, r), (8, 2)), ((a_stride_00, s), (a_stride_10, a_stride_11)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                               predicates:
                                                 8 * r == 64
                                                 a_stride_00 == (8, 16, 32, 64)[a_swizzle]
                                                 a_stride_10 == 1
                                                 (a_swizzle == 0) | (a_stride_11 == 8)
                                                 a_stride_11 % 8 == 0
                                                 k0 % 16 == 0
                                                 k0 >= 0
                                                 k0 < (8, 16, 32, 64)[a_swizzle]
                                 predicates:
                                   s % 8 == 0
       dtype=f16, form=RS      shape=(64, 16) dtype=f16 storage=rmem, held in 1 arrangement:
                                 ShardLayout(Layout((4, 2, 8, 2, 4, 2), (16, 8, 1, 512, 128, 64)), (S(0), S(2), S(4)), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
       dtype=fp8e4m3, form=SS  shape=(64, 32) dtype=fp8e4m3 storage=smem, held in 1 arrangement:
                                 a_major=K  ShardLayout(ComposedLayout(Swizzle(a_swizzle, 4, 3), k0, Layout(((8, r), (16, 2)), ((a_stride_00, s), (a_stride_10, a_stride_11)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                 predicates:
                                   8 * r == 64
                                   a_stride_00 == (16, 32, 64, 128)[a_swizzle]
                                   a_stride_10 == 1
                                   (a_swizzle == 0) | (a_stride_11 == 16)
                                   a_stride_11 % 16 == 0
                                   s % 16 == 0
                                   k0 % 32 == 0
                                   k0 >= 0
                                   k0 < (16, 32, 64, 128)[a_swizzle]
       dtype=fp8e4m3, form=RS  shape=(64, 32) dtype=fp8e4m3 storage=rmem, held in 1 arrangement:
                                 ShardLayout(Layout((4, 2, 8, 2, 4, 4), (16, 8, 1, 1024, 256, 64)), (S(0), S(2), S(4)), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
         predicates:
           each top-level mode has no backward step
           each top-level mode reaches each of its own slots exactly once
           p0 % 128 == 0
    B  dtype=bf16     shape=(16, n) dtype=bf16 storage=smem, held in 2 arrangements:
                        b_major=MN  ShardLayout(ComposedLayout(Swizzle(b_swizzle, 4, 3), 0, Layout(((8, 2), (Wb, rb)), ((b_stride_10, b_swizzle=0  lb
                                    b_swizzle=1  sb
                                    b_swizzle=2  sb
                                    b_swizzle=3  sb), (b_stride_00, b_swizzle=0  sb
                                    b_swizzle=1  lb
                                    b_swizzle=2  lb
                                    b_swizzle=3  lb)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                      predicates:
                                        Wb == (8, 16, 32, 64)[b_swizzle]
                                        rb * Wb == n
                                        b_stride_00 == 1
                                        b_stride_10 == Wb
                                        lb % 8 == 0
                        b_major=K   ShardLayout(ComposedLayout(Swizzle(b_swizzle, 4, 3), k0b, Layout(((8, 2), (8, rb)), ((b_stride_10, b_stride_11), (b_stride_00, sb)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                      predicates:
                                        8 * rb == n
                                        b_stride_00 == (8, 16, 32, 64)[b_swizzle]
                                        b_stride_10 == 1
                                        (b_swizzle == 0) | (b_stride_11 == 8)
                                        b_stride_11 % 8 == 0
                                        k0b % 16 == 0
                                        k0b >= 0
                                        k0b < (8, 16, 32, 64)[b_swizzle]
                        predicates:
                          sb % 8 == 0
       dtype=f16      shape=(16, n) dtype=f16 storage=smem, held in 2 arrangements:
                        b_major=MN  ShardLayout(ComposedLayout(Swizzle(b_swizzle, 4, 3), 0, Layout(((8, 2), (Wb, rb)), ((b_stride_10, b_swizzle=0  lb
                                    b_swizzle=1  sb
                                    b_swizzle=2  sb
                                    b_swizzle=3  sb), (b_stride_00, b_swizzle=0  sb
                                    b_swizzle=1  lb
                                    b_swizzle=2  lb
                                    b_swizzle=3  lb)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                      predicates:
                                        Wb == (8, 16, 32, 64)[b_swizzle]
                                        rb * Wb == n
                                        b_stride_00 == 1
                                        b_stride_10 == Wb
                                        lb % 8 == 0
                        b_major=K   ShardLayout(ComposedLayout(Swizzle(b_swizzle, 4, 3), k0b, Layout(((8, 2), (8, rb)), ((b_stride_10, b_stride_11), (b_stride_00, sb)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                                      predicates:
                                        8 * rb == n
                                        b_stride_00 == (8, 16, 32, 64)[b_swizzle]
                                        b_stride_10 == 1
                                        (b_swizzle == 0) | (b_stride_11 == 8)
                                        b_stride_11 % 8 == 0
                                        k0b % 16 == 0
                                        k0b >= 0
                                        k0b < (8, 16, 32, 64)[b_swizzle]
                        predicates:
                          sb % 8 == 0
       dtype=fp8e4m3  shape=(32, n) dtype=fp8e4m3 storage=smem, held in 1 arrangement:
                        b_major=K  ShardLayout(ComposedLayout(Swizzle(b_swizzle, 4, 3), k0b, Layout(((16, 2), (8, rb)), ((b_stride_10, b_stride_11), (b_stride_00, sb)))), (B(), B(), B()), Mesh(('thread',), ComposedLayout(None, p0, Layout(((4, 8, 4),), ((32, 4, 1),)))))
                        predicates:
                          8 * rb == n
                          b_stride_00 == (16, 32, 64, 128)[b_swizzle]
                          b_stride_10 == 1
                          (b_swizzle == 0) | (b_stride_11 == 16)
                          b_stride_11 % 16 == 0
                          sb % 16 == 0
                          k0b % 32 == 0
                          k0b >= 0
                          k0b < (16, 32, 64, 128)[b_swizzle]
         predicates:
           each top-level mode has no backward step
           each top-level mode reaches each of its own slots exactly once
           p0 % 128 == 0
