[Feature]: Intra-kernel profiling

[benenzhu]

Suggestion Description

Intra-kenrel profiling is really useful for communications kernels.
Also it is useful for compute-bound kernels too, to identify the XCD-imbanlance, and tail effect of some grids.

Some reference for other dsl.

1. cutedsl: https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/cute_dsl_general/iket_profiling.html
2. tilelang: [Feature] Support iket profiler for CUDA backend tile-ai/tilelang#2515
3. triton: https://github.com/triton-lang/triton/blob/main/third_party/proton/tutorials/intra_kernel/example_dsl.py
4. flashinfer: feat: flashinfer intra-kernel profiler flashinfer-ai/flashinfer#913

Also here is a blog for how to use it to identify uneven work distribution in MI300x intra-node all2all implementations:
https://gau-nernst.github.io/amd-a2a/#intra-kernel-profiling

Operating System

No response

GPU

No response

ROCm Component

No response

[LeiWang1999]

awesome, and welcome to contribute it into tilelang :)

[coderfeli]

@benenzhu thanks. make sense, we need this feature.