tree 652158ecdd8babef7457a0a38767f7833d632e66
parent 0e90e29fd0ba91100df0a704e9999fe82060c5d6
author Trevor Morris <tmorris@nvidia.com> 1684345664 -0700
committer GitHub <noreply@github.com> 1684345664 -0700
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsBcBAABCAAQBQJkZRNACRBK7hj4Ov3rIwAA/uoIAFEkbG2JzTDvU6dRo5wvIoTc
 M1c/4B1JME5QepzysAurfstNR1zyasSqpHKCufjDJGhmoDYPG/3AA8lZypS+zktQ
 Fv3rwmH+ELZgobKSvpAZ7X+JYKlkPY7v6nhSqj6fVjgFKpJZZ1hDnFOGJnYBW+YJ
 JoGi/5Nq7Y7ovWxu04E7Xr+DRP78S9l0YY3AtuF3bcLHIqEvlJlL2JfuabmRN6eR
 MKaAflS1Hwt0Eix6igStXmdzA+h7hbAJ0M5eknPpju2OgqCer9aHPRf7oNzrpEjO
 FordM59RPDHRW0moW2LEn5n5vKYhoQBA/Ah+fW95qZw2LOYzqgD7Ft0rrafzlF4=
 =SAfU
 -----END PGP SIGNATURE-----
 

Support mhlo.collective_permute with NCCL (#13502)

Adds support for `mhlo.collective_permute`.

During lowering, the `source_target_pairs` are converted into a table of
send/recv ids which can be indexed by the local rank. The send/recv id
pair is stored in the `param` field of HAL `CollectiveKind::SendRecv`.
The send and receive ids are packed as 16 bit integers into the 32 bit
param.

Fixes https://github.com/openxla/iree/issues/13100