Volta tensor op layout problem

[FindDefinition]

There is a strange line in volta smem iterator:

AccessType *access_ptr = pointer_[(s & 1) ^ ((s >> 1) & 1)];

It makes the pointers in right part of smem swapped, and recovered in

// the ' ^ ((quad & 0x4) >> 2))' swap location back
access_contiguous = ((quad & 0x4) << 1) + ((lane_in_quad) << 1) +
                      ((quad & 0x1) ^ ((quad & 0x4) >> 2));

Why should we do that? If I remove this logic, the result is correct and the speed looks same (tested in RTX2060)

[hwu36]

Which line of code? Looks like shared memory swizzling code.

[FindDefinition]

https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/transform/threadblock/regular_tile_iterator_tensor_op_sm70.h#L1193 and https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/gemm/warp/mma_tensor_op_tile_iterator_sm70.h#L1612

[hwu36]

also, what is your tile size? It maybe irrelevant to your tile size.

[FindDefinition]

using ShapeMMAThreadBlock =
    cutlass::gemm::GemmShape<128, 128, 32>; 
using ShapeMMAWarp = cutlass::gemm::GemmShape<64, 64, 32>; 
using ShapeMMAOp = cutlass::gemm::GemmShape<8, 8, 4>; 

default settings from cutlass example.

[hwu36]

TN layout, right?

[FindDefinition]

no, both are rowmajor. A use crosswise layout and B use congruous layout.

[hwu36]

how do you get rid of the logic. i tried, but get wrong results.

[FindDefinition]

1. change this line to AccessType *access_ptr = pointer_[(s & 1)];
2. change this line to

access_contiguous = ((quad & 0x4) << 1) + ((lane_in_quad) << 1) +
                          (quad & 0x1);

this change only works for rowmajor A and rowmajor B.

[hwu36]

I verified that your change can work with row x row, threadblock: 128x128, warp: 64x64.

However, when I change to row x row, threadblock: 64x128, warp: 32x64. It failed. I guess this tile size needs the swap

[FindDefinition]

actually there is another strange swap in VoltaTensorOpMultiplicandCrosswise. you need to change that line to

    int permuted_vec_contiguous =
        (vec_strided_idx & (~0xF)) + (vec_strided_idx & 0x3) * 4 +
        (((vec_strided_idx >> 2)) & 0x3);

to make 64x128x32 shape work.

[hwu36]

I made the change to the layout, too. 32x64 warp size can pass now. Then, I ran the warp level gemm unit tests and this one failed.

[  FAILED  ] 1 test, listed below:
[  FAILED  ] SM70_warp_gemm_tensor_op_crosswise.64x64x32_64x64x32_16x16x4