Fix tracing crashes for yolos, wav2vec2, hubert and convbert (aws-neuron-sdk#1265) - #1119
Open
jimburtoft wants to merge 1 commit into
Open
Conversation
These four models terminate the `torch_neuronx.trace` process on Neuron SDK >= 2.27, which is what currently blocks upgrading past 2.26. Root cause is the same for all of them, and it is not the individual ops. An `aten` op with no XLA lowering falls back to CPU via `at::native::cpu_fallback`, which calls `XLANativeFunctions::_to_cpu()` on the operands. That transfer only works if the operand was already materialized into a PjRt buffer: - operand derived from a model input -> a buffer exists -> fine - operand derived from a parameter/buffer -> torch-xla still holds an un-materialized placeholder with a null buffer -> `Check failed: pjrt_data->buffer != nullptr` -> the process dies with SIGABRT or SIGSEGV So a model only breaks when it applies an unlowered op to a *weight* instead of to an activation: wav2vec2 aten::_weight_norm_interface weight_norm parametrization hubert aten::_weight_norm_interface weight_norm parametrization yolos aten::upsample_bicubic2d position_embeddings convbert aten::im2col conv-weight-derived tensor Verifying this: the same ops at the same shapes trace fine when fed a model input, and fail when fed a parameter or a registered buffer. Note the crash classes reported upstream are misleading: the segfault and the "PjRt buffer is null" RuntimeError are one bug surfacing two ways depending on timing, and wav2vec2 was filed as a segfault but actually aborts. This adds `neuron_trace_patches`, applied from `export_neuronx`, which rewrites each call site using ops that do have XLA lowerings: - weight_norm: fold the parametrization so `w = g * v / ||v||` is evaluated once instead of per forward. Applied to every model, since any architecture using weight normalization hits this and folding is a no-op otherwise. - yolos: the bicubic interpolation depends only on a parameter and on a fixed img_size, so it is constant for a traced model. Snapshot the embeddings before tracing, then precompute and cache. The snapshot matters: reading the live tensor during tracing synchronizes an XLA tensor whose computation is still in flight and fails with `Check failed: handle->HasValue()`. - convbert: reimplement its narrow `unfold` case with pad/slice/stack/reshape, and raise on argument combinations it does not cover rather than silently computing something else. All three are numerically equivalent, verified against CPU on both the tiny-random fixtures and the real checkpoints (facebook/wav2vec2-base-960h, facebook/hubert-base-ls960, YituTech/conv-bert-base, hustvl/yolos-tiny): all four go from crash to cos_sim 1.0, with the patched-vs-unpatched CPU difference being exactly 0.0. Holds with dynamic_batch_size=True as well. Measured with tests/exporters/test_transformers.py on SDK 2.31 / inf2.8xlarge, one pytest process per model so a crash is attributable: before: yolos rc=139, convbert rc=139, wav2vec2 rc=134, hubert rc=134, cvt ok after: 18 passed, 36 skipped cvt was reported in the same issue but no longer reproduces on 2.31, so it needs no patch and could be un-skipped separately. Also relaxes requires-python to <3.13. The previous <3.12 bound reflected the Python 3.10 venvs on older Neuron DLAMIs; the SDK 2.31 DLAMI ships Python 3.12.3, where the package installs and the suites above pass. The underlying issue is still worth fixing in torch-xla: _to_cpu() could materialize the placeholder instead of asserting on it, or at minimum raise something catchable from Python rather than killing the process. Ref: aws-neuron/aws-neuron-sdk#1265
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Fixes the tracing crashes tracked in aws-neuron/aws-neuron-sdk#1265, which are what currently pin this project below Neuron SDK 2.27.
What was happening
yolos,wav2vec2,hubertandconvbertdon't raise during export — they kill the Python process, so there's nothing to catch.All four are the same bug, and it isn't the individual ops. An
atenop with no XLA lowering falls back to CPU throughat::native::cpu_fallback, which callsXLANativeFunctions::_to_cpu()on the operands. That transfer only works if the operand has already been materialized into a PjRt buffer:nullptr→Check failed: pjrt_data->buffer != nullptr→SIGABRTorSIGSEGVSo a model only breaks when it applies an unlowered op to a weight rather than to an activation:
wav2vec2aten::_weight_norm_interfaceweight_normparametrization onpos_conv_embed.convhubertaten::_weight_norm_interfaceyolosaten::upsample_bicubic2dposition_embeddingsconvbertaten::im2colThe check: the same ops at the same shapes trace fine when fed a model input, and crash when fed a parameter or a registered buffer. Only provenance changes.
Two things worth correcting from the issue as filed:
RuntimeError: PjRt buffer is nullare the same null-buffer dereference surfacing differently depending on timing.wav2vec2was filed under segfault but actually aborts.cvtno longer reproduces on 2.31. It passes untouched, so it needs no patch and could be un-skipped in a separate PR.What this PR does
Adds
optimum/exporters/neuron/neuron_trace_patches.py, applied fromexport_neuronx, which rewrites each call site using ops that do have XLA lowerings so no CPU fallback is attempted:weight_norm— fold the parametrization sow = g * v / ||v||is evaluated once eagerly instead of on every forward. Applied to every model rather than just these two: any architecture using weight normalization (speech encoders, vocoders, some GANs) hits this identically, and folding is a no-op otherwise.yolos— the bicubic interpolation depends only on a parameter and onimg_size, which is fixed for a traced model, so it's constant with respect to the graph. Snapshot the embeddings before tracing, then precompute and cache. The snapshot matters: reading the live tensor mid-trace synchronizes an XLA tensor whose computation is still in flight and fails withCheck failed: handle->HasValue().convbert— reimplement its narrowunfoldcase with pad/slice/stack/reshape. It raises on argument combinations it doesn't cover rather than silently computing something else.All are monkey-patches applied at runtime, so no vendored
transformerschanges.Verification
tests/exporters/test_transformers.pyon SDK 2.31 / inf2.8xlarge, one pytest process per model so a crash is attributable rather than taking down the run:yolosrc=139(SIGSEGV)wav2vec2rc=134(SIGABRT)convbertrc=139(SIGSEGV)hubertrc=134(SIGABRT)cvtTogether: 18 passed, 36 skipped (previously 2 failed / 16 passed, and the run aborted partway when invoked as one process).
Numerical equivalence, against CPU, on the tiny-random fixtures and the real checkpoints:
facebook/wav2vec2-base-960hfacebook/hubert-base-ls960YituTech/conv-bert-basehustvl/yolos-tinyThe last column is the point: the patches are numerically inert. Residual
cos_simdeviation is ordinary FP32 Neuron-vs-CPU noise. Also verified withdynamic_batch_size=True, sincetest_compare_to_transformers_dyn_bswas one of the failing tests.New tests in
tests/exporters/test_neuron_trace_patches.py(13 cases, CPU-only, no Neuron needed, ~13s):unfold_via_slicesexactness vsnn.functional.unfold, one case per clause of its argument guard, weight-norm folding equivalence, and output equivalence for all four models.requires-pythonAlso relaxes
requires-pythonto<3.13. The previous<3.12bound reflected the Python 3.10 venvs on older Neuron DLAMIs, but the SDK 2.31 DLAMI ships Python 3.12.3 — the package wouldn't install there at all. Everything above was run on 3.12.3 in the stockaws_neuronx_venv_pytorch_2_9venv. Happy to split this into its own PR if you'd rather keep it separate.Still worth fixing upstream
These are workarounds; the real issue is in torch-xla.
_to_cpu()could materialize the placeholder instead of asserting on it — parameters are known constants at trace time. Failing that, replacing theCHECKwith a thrown exception would at least make it catchable from Python instead of killing the process. I've noted this on the Neuron issue.Environment
SDK 2.31 (DLAMI 20260708), inf2.8xlarge, torch-neuronx
2.9.0.2.15.32035, neuronx-cc2.26.6360.0, torch2.9.1, transformers4.57.6, Python3.12.3.