e.g. 4行, 8列, 2B数据
index: 0b rr ccc b
rr表示行号所在bits
ccc表示列号所在bits
b表示数据组所需bits

swizzle就相当于把ccc替换成ccc ^ rr

// A generic Swizzle functor
/* 0bxxxxxxxxxxxxxxxYYYxxxxxxxZZZxxxx
 *                               ^--^ MBase is the number of least-sig bits to keep constant
 *                  ^-^       ^-^     BBits is the number of bits in the mask
 *                    ^---------^     SShift is the distance to shift the YYY mask
 *                                       (pos shifts YYY to the right, neg shifts YYY to the left)
 *
 * e.g. Given
 * 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx
 * the result is
 * 0bxxxxxxxxxxxxxxxxYYxxxxxxxxxAAxxx where AA = ZZ xor YY
 */
 template <int BBits, int MBase, int SShift = BBits>
struct Swizzle
{
  static constexpr int num_bits = BBits;
  static constexpr int num_base = MBase;
  static constexpr int num_shft = SShift;

  static_assert(num_base >= 0,             "MBase must be positive.");
  static_assert(num_bits >= 0,             "BBits must be positive.");
  static_assert(abs(num_shft) >= num_bits, "abs(SShift) must be more than BBits.");

  // using 'int' type here to avoid unintentially casting to unsigned... unsure.
  // NOTE: 数据位宽
  using bit_msk = cute::constant<int, (1 << num_bits) - 1>;
  // NOTE: "行号"位
  using yyy_msk = cute::constant<int, bit_msk{} << (num_base + max(0,num_shft))>;
  // NOTE: "列号"位
  using zzz_msk = cute::constant<int, bit_msk{} << (num_base - min(0,num_shft))>;
  using msk_sft = cute::constant<int, num_shft>;

  static constexpr uint32_t swizzle_code = uint32_t(yyy_msk::value | zzz_msk::value);

  template <class Offset>
  CUTE_HOST_DEVICE constexpr static
  auto
  apply(Offset const& offset)
  {
    // NOTE:
    // 1. shiftr(offset & yyy_msk{}, msk_sft{})
    //    shiftr表示右移
    //    offset & yyy_msk{} >> shift
    //    只保留yyy_bis, 其他bit清0
    //    yyy_bis >> shift
    //    000000yyy000 >> shift
    // 2. offset ^ (yyy_bis >> shift)
    //  0bxxxxxxxxxxxxxxxxYYxxxxxxxxxZZxxx
    //       xor
    //  0b000000000000000000000000000YY000
    return offset ^ shiftr(offset & yyy_msk{}, msk_sft{});   // ZZZ ^= YYY
  }

  template <class Offset>
  CUTE_HOST_DEVICE constexpr
  auto
  operator()(Offset const& offset) const
  {
    return apply(offset);
  }
};