Skip to content

[Neuron] Enable torch.compile compatibility with Neuron device#13485

Merged
yiyixuxu merged 38 commits into
huggingface:mainfrom
JingyaHuang:neuron-torch-comppile
Jun 18, 2026
Merged

[Neuron] Enable torch.compile compatibility with Neuron device#13485
yiyixuxu merged 38 commits into
huggingface:mainfrom
JingyaHuang:neuron-torch-comppile

Conversation

@JingyaHuang

@JingyaHuang JingyaHuang commented Apr 15, 2026

Copy link
Copy Markdown
Contributor

What does this PR do?

Following #13289

Adding torch.compile support for Neuron to diffusers.

Neuron supports torch.compile with the following mechanism to improve performance:

torch.compile(model, backend="neuron", fullgraph=True)                                                                                                 
           │                                                
           ▼                                                                                                                                             
     TorchDynamo                   traces Python → FX graph (aten ops)
           │                                                                                                                                             
           ▼                                                
     AOT Autograd                  lifts params as inputs, decomposes ops                                                                                
           │                                                                                                                                             
           ▼
     Neuron backend                Neuron-specific decompositions                                                                                        
     (torch_neuronx)               (int64→int32, SDPA lowering, NKI passes)
           │  StableHLO                                                                                                                                  
           ▼                                                                                                                                             
     neuronx-cc                    compiles HLO → NEFF binary                                                                                            
                                   allocates on-chip SRAM (SB, 28 MB/core)                                                                               
           │  NEFF                                                                                                                                       
           ▼                                                                                                                                             
     NEFF cache                    keyed by StableHLO hash                                                                                               
     (TORCH_NEURONX_NEFF_CACHE_DIR)  miss → recompile; hit → load instantly                                                                              
           │                                                                                                                                             
           ▼                                                                                                                                             
     Executor                      preprocesses inputs (contiguous, int64→int32)                                                                         
     (executor.py)                 calls _C.execute_compiled_graph() 

This PR integrates torch.compile on Neuron seamlessly into Diffusers via the following changes:

  • In src/diffusers/models/transformers/transformer_flux2.py

    • Flux2PosEmbed.forward() — add Neuron to the float32 RoPE guard (Neuron doesn't support float64, same pattern as MPS/NPU)
    • Flux2SwiGLU.forward() — replace x.chunk(2, dim=-1) with explicit x[..., :half] / x[..., half:] slices (equivalent on all backends; avoids a gather op in the Neuron IR that causes NCC_ITEN404 DataLocalityOpt compilation failure under torch.compile)
  • In src/diffusers/models/unets/unet_2d_condition.py

    • get_time_embed() — use int32 timestep dtype on Neuron (Neuron doesn't support int64)
  • In src/diffusers/pipelines/stable_diffusion_xl/pipeline_stable_diffusion_xl.py

    • Keep timesteps on CPU for Neuron (same pattern as XLA);
    • cast to float32 before moving to device;
    • run scheduler step on CPU output
  • In src/diffusers/utils/import_utils.py + torch_utils.py

    • Register the neuron backend: BACKEND_SYNCHRONIZE, BACKEND_SUPPORTS_TRAINING, BACKEND_DEVICE_COUNT, randn_tensor forces CPU random generation then .to(device) (Neuron has no native RNG)
  • a test for pixart

Validation

So far we validated the following models, the idea is to manually validated a part of representative models, and ensure the model coverage with pur CIs afterward.

  • pixart
  • sdxl
  • flux2-klein-4B (able to run within one NeuronCore w/o. tp)

Who can review?

Anyone in the community is free to review the PR once the tests have passed. Feel free to tag
members/contributors who may be interested in your PR.

Caveats

Somes fixes need to be upsteamed to TorchNeuron, this PR won't be stable with potential changes in the SDK

@HuggingFaceDocBuilderDev

Copy link
Copy Markdown

The docs for this PR live here. All of your documentation changes will be reflected on that endpoint. The docs are available until 30 days after the last update.

@github-actions github-actions Bot added size/L PR with diff > 200 LOC and removed size/L PR with diff > 200 LOC labels Apr 15, 2026
@JingyaHuang JingyaHuang changed the title [Neuron] Enable torch.comppile compatibility with Neuron device [Neuron] Enable torch.compile compatibility with Neuron device Apr 21, 2026
@github-actions github-actions Bot removed the size/L PR with diff > 200 LOC label Apr 22, 2026
@github-actions github-actions Bot added size/M PR with diff < 200 LOC and removed size/L PR with diff > 200 LOC labels Jun 3, 2026
@github-actions github-actions Bot added size/L PR with diff > 200 LOC and removed size/M PR with diff < 200 LOC labels Jun 3, 2026
@github-actions github-actions Bot added size/M PR with diff < 200 LOC and removed size/L PR with diff > 200 LOC labels Jun 3, 2026
@JingyaHuang JingyaHuang requested a review from sayakpaul June 4, 2026 16:52
# We set the index here to remove DtoH sync, helpful especially during compilation.
# Check out more details here: https://github.com/huggingface/diffusers/pull/11696
self.scheduler.set_begin_index(0)
if hasattr(self.scheduler, "set_begin_index"):

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This seems like a bigger change that should be propagated to other pipelines that use self.scheduler.set_begin_index(0)?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yes, only a part of self.scheduler have set_begin_index, so it will be good to have the check, I only added for the pipelines related to the PR, do you want me to propagate to others? Or I could open a separate PR for this.

Comment on lines +910 to +911
if hasattr(self.scheduler, "set_begin_index"):
self.scheduler.set_begin_index(0)

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If the pipeline initially didn't have it, maybe we don't need it?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The change was suggested in the initial PR for supporting neuron: #13289 (comment).

the line is needed for the compiled path otherwise the first scheduler.step calls index_for_timestep, which does a comparison/nonzero and forces a DtoH sync inside the compiled denoising loop (same motivation as #11696).

@sayakpaul sayakpaul left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Left some further comments.

@JingyaHuang JingyaHuang left a comment

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks @sayakpaul for the review! I addressed the comments you made.

Comment on lines -294 to +295
x1, x2 = x.chunk(2, dim=-1)
x = self.gate_fn(x1) * x2
half = x.shape[-1] // 2
x = self.gate_fn(x[..., :half]) * x[..., half:]

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yeah, here is totally a fix for neuron compilation. neuron compiler lower chunk to split/gather and previously introduced a compilation failure, now it's fixed by neuron's new beta drop. But the explicit static slices x[..., :half] / x[..., half:] lower to clean slice ops, and the neuron compiler compiles it faster.

Comment on lines +910 to +911
if hasattr(self.scheduler, "set_begin_index"):
self.scheduler.set_begin_index(0)

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The change was suggested in the initial PR for supporting neuron: #13289 (comment).

the line is needed for the compiled path otherwise the first scheduler.step calls index_for_timestep, which does a comparison/nonzero and forces a DtoH sync inside the compiled denoising loop (same motivation as #11696).

# We set the index here to remove DtoH sync, helpful especially during compilation.
# Check out more details here: https://github.com/huggingface/diffusers/pull/11696
self.scheduler.set_begin_index(0)
if hasattr(self.scheduler, "set_begin_index"):

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

yes, only a part of self.scheduler have set_begin_index, so it will be good to have the check, I only added for the pipelines related to the PR, do you want me to propagate to others? Or I could open a separate PR for this.

def test_flux2_klein_neuron_compile_128(self):
import torch_neuronx # noqa: F401 — registers torch.neuron
from torch_neuronx.neuron_dynamo_backend import set_model_name
from transformers.utils.output_capturing import install_all_output_capturing_hooks

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It was added for the compilation of the text encoder, which is not a component we compile in general, removed it!

@yiyixuxu yiyixuxu merged commit 236e5dd into huggingface:main Jun 18, 2026
21 of 22 checks passed
@yiyixuxu

Copy link
Copy Markdown
Collaborator

thanks a ton for working on this @JingyaHuang !

@JingyaHuang JingyaHuang deleted the neuron-torch-comppile branch June 18, 2026 17:25
DN6 pushed a commit that referenced this pull request Jul 1, 2026
)

* draft:add neuron as a legit backend

* feat: neuron-specific changes in the pipeline

* tests: eager tests

* fix: style

* fix:apr_02 beta

* fix:pixart

* fix: rewrite flux swiglu activation to avoid gather op in neuron IR

* test: pixart compile mode on neuron

* cleanup & fix style

* test: flux2 klein and sdxl-turbo

* review: address some comments

* revert:put back expected slices

* cleanup:move the ts index ahead of denoising + cleanup tests

* fix: style

* fix style

* review:remove the hook

---------

Co-authored-by: Sayak Paul <spsayakpaul@gmail.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants