[EVT] Add support for Row/Col broadcast PtrArray

[jwfromm]

To enable FP8 grouped gemm with rowwise scaling in the epilogue, we need to be able to provide a list of pointers to the scales for each group. This PR extends Sm90Row/ColBroadcast to support PTRArray to handle this case. Now if the ElementInput type is specified as a pointer, the corresponding input is an array of pointers, enabling grouped gemm. For example in our case we define EVT nodes like this:

  using XScale = cutlass::epilogue::fusion::Sm90ColBroadcast<
      0,
      TileShape,
      ElementComputeEpilogue*, // Indicate input is array of pointers.
      ElementComputeEpilogue,
      cute::Stride<cute::Int<1>, cute::Int<0>, cute::Int<0>>>;

  using WScale = cutlass::epilogue::fusion::Sm90RowBroadcast<
      0,
      TileShape,
      ElementComputeEpilogue*, // Indicate input is array of pointers.
      ElementComputeEpilogue,
      cute::Stride<cute::Int<0>, cute::Int<1>, cute::Int<0>>>;

[jiawenliu64]

@ANIKET-SHIVAM @hwu36 Do you have a timeline to review this? We need this feature enabled on cutlass ASAP to unblock our usecases at Meta, e.g., pytorch/FBGEMM#3560

[jwfromm]

Thanks for taking a look @Skylion007, I've incorporated your feedback if you'd like to check again to make sure this all looks good.

[Skylion007]

LGTM

[jiawenliu64]

Thanks! Can you merge this to cutlass to unblock e.g., pytorch/FBGEMM#3560?

[Skylion007]

Need a nvidia employee to to do that. @eqy might know who to ping.

[jiawenliu64]

cc. @hwu36

[hwu36]

We are working on this.

[jwfromm]

@ANIKET-SHIVAM I've refactored this PR so that the functionality is built into the existing Row/Col EVT nodes. Can you take another look?