Get remote tensors inside `@helion.kernel` by kwen2501 · Pull Request #1122 · pytorch/helion

kwen2501 · 2025-11-13T03:09:49Z

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 · 2025-11-13T06:15:57Z

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

yf225 · 2025-11-13T07:45:50Z

@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

examples/all_reduce.py

yf225 · 2025-11-13T19:32:03Z

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

kwen2501 · 2025-11-14T02:00:31Z

@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 · 2025-11-14T04:17:21Z

@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 !

To support use case in pytorch/helion#1122, i.e. ``` @helion.kernel def foo( x: Tensor, group_name: str ): x_remotes = torch.ops.symm_mem.get_remote_tensors(x, group_name) for t in x_remotes: ... ```` Helion uses fake tensor to trace a program, thus we cannot use the following code in a Helion function: ``` hdl = rendezvous(tensor) remote_tensors = tuple( hdl.get_remote_tensor(peer, ...) for peer in range(world_size) ) ``` The reason is that when `tensor` is fake, the returned `hdl` is None, thus any subsequent call on it will fail. This PR wraps the above functionality as an op: ``` lib.define("get_remote_tensors(Tensor x, str group_name) -> Tensor[]") ``` so that things like `hdl` is not exposed to Helion. The op also provides a `meta` implementation so that Helion can trace it without actually running the rendezvous. Pull Request resolved: #167779 Approved by: https://github.com/yf225

To support use case in pytorch/helion#1122, i.e. ``` @helion.kernel def foo( x: Tensor, group_name: str ): x_remotes = torch.ops.symm_mem.get_remote_tensors(x, group_name) for t in x_remotes: ... ```` Helion uses fake tensor to trace a program, thus we cannot use the following code in a Helion function: ``` hdl = rendezvous(tensor) remote_tensors = tuple( hdl.get_remote_tensor(peer, ...) for peer in range(world_size) ) ``` The reason is that when `tensor` is fake, the returned `hdl` is None, thus any subsequent call on it will fail. This PR wraps the above functionality as an op: ``` lib.define("get_remote_tensors(Tensor x, str group_name) -> Tensor[]") ``` so that things like `hdl` is not exposed to Helion. The op also provides a `meta` implementation so that Helion can trace it without actually running the rendezvous. Pull Request resolved: pytorch#167779 Approved by: https://github.com/yf225

meta-cla · 2025-11-22T02:14:48Z

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!

meta-cla bot added the CLA Signed This label is managed by the Meta Open Source bot. label Nov 13, 2025

yf225 reviewed Nov 13, 2025

View reviewed changes

kwen2501 force-pushed the get_remote_tensors branch from 9ec6b6e to 7fda194 Compare November 13, 2025 18:21

yf225 reviewed Nov 13, 2025

View reviewed changes

examples/all_reduce.py Show resolved Hide resolved

yf225 force-pushed the get_remote_tensors branch from 7fda194 to 47f5294 Compare November 13, 2025 19:32

kwen2501 mentioned this pull request Nov 14, 2025

[SymmMem] op to get remote tensors pytorch/pytorch#167779

Closed

yf225 approved these changes Nov 14, 2025

View reviewed changes

jansel approved these changes Nov 16, 2025

View reviewed changes

kwen2501 added 4 commits November 17, 2025 11:53

Get remote tensors inside Helion kernel

fe310bb

Update tests

66809c8

fix lint

fd8dc16

Add todo

1db590a

kwen2501 force-pushed the get_remote_tensors branch from f342333 to 1db590a Compare November 17, 2025 19:53

yf225 added 2 commits November 22, 2025 18:34

Merge branch 'main' into get_remote_tensors

f022fd8

Merge branch 'main' into get_remote_tensors

c4f37e2

yf225 merged commit 198f1cb into pytorch:main Nov 24, 2025
16 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Get remote tensors inside `@helion.kernel`#1122

Get remote tensors inside `@helion.kernel`#1122
yf225 merged 6 commits intopytorch:mainfrom
kwen2501:get_remote_tensors

kwen2501 commented Nov 13, 2025 •

edited

Loading

Uh oh!

yf225 left a comment

Uh oh!

yf225 commented Nov 13, 2025 •

edited

Loading

Uh oh!

yf225 commented Nov 13, 2025

Uh oh!

Uh oh!

yf225 commented Nov 13, 2025

Uh oh!

kwen2501 commented Nov 14, 2025

Uh oh!

yf225 commented Nov 14, 2025

Uh oh!

yf225 left a comment

Uh oh!

meta-cla bot commented Nov 22, 2025

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

kwen2501 commented Nov 13, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

yf225 left a comment

Choose a reason for hiding this comment

Uh oh!

yf225 commented Nov 13, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

yf225 commented Nov 13, 2025

Uh oh!

Uh oh!

yf225 commented Nov 13, 2025

Uh oh!

kwen2501 commented Nov 14, 2025

Uh oh!

yf225 commented Nov 14, 2025

Uh oh!

yf225 left a comment

Choose a reason for hiding this comment

Uh oh!

meta-cla bot commented Nov 22, 2025

Process

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

kwen2501 commented Nov 13, 2025 •

edited

Loading

yf225 commented Nov 13, 2025 •

edited

Loading