[TIRx][CUDA] Support predicated gemm_async

[jinhongyii]

Motivation and context

Blackwell tcgen05.mma is issued by a single thread, while one Tx.gemm_async tile call may expand into multiple dense or block-scaled MMA instructions over its M, N, and K iterations. Persistent kernels sometimes need to keep the surrounding state machine warp-uniform, so they compute a one-lane predicate and apply it at the instruction level:

mma_issue = T.ptx.elect_sync()
Tx.gemm_async(
    C,
    A,
    B,
    dispatch="tcgen05",
    pred=mma_issue,
)
mma_bar.arrive(stage, pred=mma_issue)

Before this PR, the pred configuration was not forwarded by the gemm_async tcgen05 dispatcher to every MMA it generated. In addition, the block-scaled MMA path and TCGen05Bar.arrive did not support forwarding the same predicate. Predicating only the matching commit is not sufficient: the active lanes may still issue all internal MMA instructions.

Changes

• Read the optional pred value from the gemm_async tile-call configuration and forward it to every generated dense or block-scaled tcgen05 MMA.
• Add optional instruction-predicate support to ptx_tcgen05_mma_block_scale and its CUDA helper generation.
• Generate a predicated PTX instruction by converting the runtime uint32 value to a PTX predicate and emitting @p_issue tcgen05.mma....
• Add pred to TCGen05Bar.arrive and forward it to both the single-CTA and multicast tcgen05.commit forms.
• Document the supported thread-scope pattern and the requirement that the MMA sequence and matching commit use the same predicate.

The existing behavior remains unchanged when pred is omitted.

Testing

• Verify predicated single-CTA and two-CTA multicast TCGen05Bar commits in CUDA codegen.
• Verify that block-scaled MMA codegen emits @p_issue tcgen05.mma....
• Verify that one dense or block-scaled tile call expands to four internal MMA calls and that all calls receive the same predicate as the matching commit.
• Run changed-files pre-commit checks.

Focused result: 5 tests passed.

[gemini-code-assist]

Caution

The consumer version of Gemini Code Assist on GitHub has been sunset. All code review activity has officially ceased.

[jinhongyii]

Closing this in favor of the existing explicit single-lane guard form (if mma_issue:). On the current grouped FP8 kernel, a like-for-like B200 A/B benchmark found no measurable regression from the explicit guard, and the generated SASS still uses uniform registers for the MMA descriptor path. The additional pred API therefore does not appear justified by this use case.