[QST] Why does slice_and_offset() (in swizzle_layout.hpp) not consider smem_ptr_flag_bits ?

What is your question?
Based on my understanding of the implementation of slice_and_offset() for swizzle layout, it tries to make the input swizzle layout decay to a normal layout. The implementation seems to assume that the input swizzle layout is position independent? But what if the Offset is smem_ptr_flag_bits? smem_ptr_flag_bits will make the swizzle layout position dependent. Would it decay to a normal layout mistakenly?

I noticed there exist slice operations for smem ptr flagged swizzle layouts in CUTLASS, which will call slice_and_offset().

E.g.
typename Params::TMA_A tma_load_a = make_tma_copy_A_sm90(
GmemTiledCopyA{},
tensor_a,
SmemLayoutA{}(,,cute::Int<0>{}),
TileShape{},
ClusterShape{});

SmemLayoutA is indeed an smem ptr flagged swizzle layout. Does it just happen not to decay to a normal layout?

Start in swizzle_layout.hpp at slice_and_offset(), then trace how smem_ptr_flag_bits is represented in the flagged swizzle layouts used by the TMA_A example. Compare the slice operations for smem pointer flagged layouts and determine whether the decay preserves position dependence; done means a confirmed explanation with a focused correction if the behavior is wrong.