Capture MVP requirements for thread groups from Hierarchical Parallelism use case

[jrhemstad]

This issue can be closed with a description of the minimal API required from new cuda:: thread groups to satisfy requirements from hierarchical parallel algorithms.

[NaderAlAwar]

The thread groups today already have most of what is required for Hierarchical Parallelism. Some additional nice-to-haves would be:

• A method to elect one thread. This could be similar to invoke_one and invoke_one_broadcast https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cooperative-groups.html#invoke-one, or maybe we can also provide an elect(), which would allow the caller to reuse the elected thread
• Some method to exchange data between threads, like a shuffle or cub::BlockExchange
• A method to broadcast the result from one thread to others in the group. I needed something like this to implement cuda::coop::reduce() over a block group, I ended up manually creating an smem variable and populating it. Would be nice to hide that.

The thread groups APIs I am currently using:

• cuda::experimental::this_block{block_hierarchy};: construct blocks that are passed to the cooperative prologue/epilogue, enabling things like cuda::coop::reduce()
• cuda::gpu_thread.extents(...)::static_extent(0): Used in cuda::coop::reduce() to get threads per block to pass to cub::BlockReduce
• cuda::gpu_thread.is_root_rank(group): Used in cuda::coop::reduce() to get the root thread which holds the result of the cub::BlockReduce
• group.sync(): needed after reduction and generally useful all around

[bernhardmgruber]

A method to elect one thread

We need this as well in case we want to replace the warpspeed squads by the new groups.

[NaderAlAwar]

Finished the PoC implementation for RMSNorm hierarchical parallelism, details here #8358 (comment). I believe the items I mentioned in my above comment are valuable (cc @davebayer), but beyond that I don't see any other requirements.

[jrhemstad]

@NaderAlAwar could you update your comment to include a list of all thread group APIs you are using? Would like to have a concrete list.

[NaderAlAwar]

@jrhemstad added to my original comment, also added broadcast as a requirement