[QST] How to print values in register fragments via DescriptorIterator?

[hyhieu]

I have a kernel that looks somewhat like this (relevant lines only):

auto gA = local_tile(
  tma_a.get_tma_tensor(shape(A)),
  Shape<BLK_M, BLK_K>{},
  make_coord(cta_x, _));
auto sA = make_tensor(make_smem_ptr(shared_storage.smem_a.data()), SmemLayoutA{});
auto tiled_mma = TiledMma{};
auto thr_mma = tiled_mma.get_slice(thr_idx);
auto tCsA = thr_mma.partition_A(sA);
auto tCrA = thr_mma.make_fragment_A(tCsA);
if (cute::thread0()) {
  print_tensor(tCrA);
}

Then the printed output looks like:

GMMA::DescriptorIterator o (_1,_1,_1):(_0,_0,_0):
GMMA::DescriptorIterator o (_1,_1):(_0,_0):
  GmmaDescriptor: 0x0000000800400040
  start_addr :  0x0040
  leading_off:  0x0040 (64)
  stride_off :  0x0008 (8)
  base_offset:  0x0
  layout_type:  0x0 (INTERLEAVE)

As such, I don't really the values in tCrA. How can I really see the values in tCrA?

Thanks!

[ccecka]

Those are the values in tCrA :-) The values of the tensor are SMEM Descriptors.

Those descriptors describe the values in tCsA, so I believe you want to print those.

[hyhieu]

Wow. I think this answer elicits a lot of my understanding of the concept of GMMA::DescriptorIterator.

Thank you, Cris! As you mentioned you were working on more documentations for CuTe, I think the topic of DescriptorIterator would is very interesting and helpful.