Why is the base_offset field of the WGMMA matrix descriptor not set according to PTX requirement?

[tianlei-zhao]

As the PTX manual states:

The value of base offset is 0 when the repeating pattern of the specified swizzling mode starts as per the below table:

Swizzling mode       Starting address of the repeating pattern
128-Byte swizzle     1024-Byte boundary
 64-Byte swizzle     512-Byte boundary
 32-Byte swizzle     256-Byte boundary

Otherwise, the base offset must be a non-zero value, computed using the following formula:
base offset = (pattern start addr >> 0x7) & 0x7

However, CUTLASS set the base offset filed as 0 constantly, even if the SMEM A/B buffer is only aligned on a 128byte boundary:

    struct TensorStorage : cute::aligned_struct<128, _0> {
      cute::array_aligned<typename TiledMma::ValTypeA, cute::cosize_v<SmemLayoutA>> smem_A;
      cute::array_aligned<typename TiledMma::ValTypeB, cute::cosize_v<SmemLayoutB>> smem_B;
    } tensors;

...

  constexpr uint8_t base_offset = 0;
  desc.bitfield.base_offset_ = base_offset;

So I'm wondering what is the purpose of the base_offset filed? Why is it not set according to the PTX manual?

Thanks

[tianlei-zhao]

@thakkarV @hwu36 Sorry to bother you guys, but this really puzzled me a long time, besides the PTX doc is vague here.