[Neuron] Enable torch.compile compatibility with Neuron device#13485
Conversation
… into add-neuron-backend
… into add-neuron-backend
|
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. |
torch.comppile compatibility with Neuron devicetorch.compile compatibility with Neuron device
| # 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"): |
There was a problem hiding this comment.
This seems like a bigger change that should be propagated to other pipelines that use self.scheduler.set_begin_index(0)?
There was a problem hiding this comment.
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.
| if hasattr(self.scheduler, "set_begin_index"): | ||
| self.scheduler.set_begin_index(0) |
There was a problem hiding this comment.
If the pipeline initially didn't have it, maybe we don't need it?
There was a problem hiding this comment.
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
left a comment
There was a problem hiding this comment.
Left some further comments.
JingyaHuang
left a comment
There was a problem hiding this comment.
Thanks @sayakpaul for the review! I addressed the comments you made.
| 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:] |
There was a problem hiding this comment.
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.
| if hasattr(self.scheduler, "set_begin_index"): | ||
| self.scheduler.set_begin_index(0) |
There was a problem hiding this comment.
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"): |
There was a problem hiding this comment.
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 |
There was a problem hiding this comment.
It was added for the compilation of the text encoder, which is not a component we compile in general, removed it!
|
thanks a ton for working on this @JingyaHuang ! |
) * 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>
What does this PR do?
Following #13289
Adding torch.compile support for Neuron to diffusers.
Neuron supports
torch.compilewith the following mechanism to improve performance:This PR integrates
torch.compileon Neuron seamlessly into Diffusers via the following changes:In
src/diffusers/models/transformers/transformer_flux2.pyFlux2PosEmbed.forward()— add Neuron to thefloat32RoPE 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 undertorch.compile)In
src/diffusers/models/unets/unet_2d_condition.pyget_time_embed()— use int32 timestep dtype on Neuron (Neuron doesn't support int64)In
src/diffusers/pipelines/stable_diffusion_xl/pipeline_stable_diffusion_xl.pyIn
src/diffusers/utils/import_utils.py + torch_utils.py.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.
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