Get remote tensors inside @helion.kernel

[kwen2501]

Instead of user passing a tuple of tensors into the kernel.

 def one_shot_all_reduce_kernel(
     signal_pad_addrs: torch.Tensor,
     local_signal_pad: torch.Tensor,
-    a_shared_tuple: tuple[torch.Tensor, ...],
+    a_shared: torch.Tensor,

We get the tuple of remote tensors by calling torch.ops.symm_mem.get_remote_tensors in the CPU part of the Helion function.

This op is yet to be upstreamed on PyTorch side. Naively, it is nothing but:

lib.define(
    "get_remote_tensors(Tensor x, str group_name) -> Tensor[]"
)

@torch.library.impl(lib, "get_remote_tensors", "CUDA")
def _get_remote_tensors_default(
    local: torch.Tensor,
    group_name: str
):
    hdl = rendezvous(local, group_name)
    return tuple(
        hdl.get_remote_tensor(peer, local.size(), local.dtype) for peer in range(hdl.world_size)
    )

@torch.library.impl(lib, "get_remote_tensors", "Meta")
def _get_remote_tensors_meta(
    local: torch.Tensor,
    group_name: str
):
    # TODO: correct world size
    world_size = torch.distributed.get_world_size()
    return tuple(
        torch.empty_like(local) for _ in range(world_size)
    )

The "Meta" impl is necessary because Helion seems to traces the function in Fake mode.

[yf225]

Thanks! I believe we need to update test_examples_dist.py as well, but Helion distributed CI currently has a bug causing the test error to not surface. I will land a PR to fix the bug and then we can rebase this PR

[yf225]

@kwen2501 I will try to see if I can get the distributed CI unit test to show up as fail on #1125

[yf225]

@kwen2501 in case adding get_remote_tensors to PyTorch core would take some time, we can also define the custom op at top of the Helion example file:

lib = torch.library.Library("symm_mem", "FRAGMENT")  # noqa: TOR901

lib.define(
    "get_remote_tensors(Tensor x, str group_name) -> Tensor[]"
)

@torch.library.impl(lib, "get_remote_tensors", "CUDA")
def _get_remote_tensors_default(
    local: torch.Tensor,
    group_name: str
):
    hdl = torch.distributed._symmetric_memory.rendezvous(local, group_name)
    return tuple(
        hdl.get_remote_tensor(peer, local.size(), local.dtype) for peer in range(hdl.world_size)
    )

@torch.library.impl(lib, "get_remote_tensors", "Meta")
def _get_remote_tensors_meta(
    local: torch.Tensor,
    group_name: str
):
    # TODO: correct world size
    world_size = torch.distributed.get_world_size()
    return (local,) * world_size

[yf225]

@kwen2501 I'll rebase this PR so that it has the distributed CI error propagation fix. Thanks!

[kwen2501]

@yf225 What torch version does CI use?
It is missing torch.ops.symm_mem.get_remote_tensors. Would it be picked up after we land it in torch nightly?

[yf225]

@yf225 What torch version does CI use? It is missing torch.ops.symm_mem.get_remote_tensors. Would it be picked up after we land it in torch nightly?

yes it uses torch nightly - should be able to pick it up very soon

[yf225]

thanks a lot @kwen2501 !

[meta-cla]

Hi @kwen2501!

Thank you for your pull request.

We require contributors to sign our Contributor License Agreement, and yours needs attention.

You currently have a record in our system, but the CLA is no longer valid, and will need to be resubmitted.

Process

In order for us to review and merge your suggested changes, please sign at https://code.facebook.com/cla. If you are contributing on behalf of someone else (eg your employer), the individual CLA may not be sufficient and your employer may need to sign the corporate CLA.

Once the CLA is signed, our tooling will perform checks and validations. Afterwards, the pull request will be tagged with CLA signed. The tagging process may take up to 1 hour after signing. Please give it that time before contacting us about it.

If you have received this in error or have any questions, please contact us at cla@meta.com. Thanks!