fix(rope): don't mark the unused batch-size param as tl.constexpr - #1351
Draft
vaibhavjindal wants to merge 2 commits into
Draft
fix(rope): don't mark the unused batch-size param as tl.constexpr#1351vaibhavjindal wants to merge 2 commits into
vaibhavjindal wants to merge 2 commits into
Conversation
`_triton_rope` declares `bs: tl.constexpr`, but `bs` is never referenced in
the kernel body -- it appears only in comments. Marking a dead parameter as
constexpr still forces Triton/Dynamo to specialize on it, and that breaks
torch.compile as soon as the batch dimension becomes dynamic: the value
arrives as a SymInt, and Dynamo trips an internal assertion while tracing
LigerRopeFunction through the autograd_function_apply HOP:
assert subgraph_vt.is_tensor() or isinstance(subgraph_vt, SymNodeVariable)
torch/_dynamo/variables/higher_order_ops.py:325
The failure only surfaced on the *second* distinct input shape, since the
first call compiles statically and the second triggers automatic dynamic
shapes -- which made it look like an upstream torch bug rather than a kernel
signature issue.
Varying seq_len alone was always fine, because `sl` is an ordinary runtime
arg; only the constexpr batch size was fatal:
before after
vary seq_len only (auto-dynamic) OK OK
vary batch only (auto-dynamic) FAIL OK
vary batch, dynamic=False OK OK
Dropping the constexpr annotation is sufficient. Qwen3-0.6B now compiles
under a plain `torch.compile(model)` across batch sizes 1/4/8 and sequence
lengths 512/1024/2048; previously that required a `dynamic=False` workaround.
Explicit `torch._dynamo.mark_dynamic` on the batch and sequence dims also
works now.
The parameter is kept (rather than removed) so the kernel's call signature
and the `q size: (bsz, ...)` documentation stay intact.
Note `dynamic=True` remains unsupported: it additionally marks head_dim and
the head counts dynamic, which cannot work while block sizing is derived from
`triton.next_power_of_2(head_dim)` at launch time. Those dims are fixed by
model architecture in practice, so this is not the case that matters.
Behaviour is unchanged on the eager path, and compiled output and gradients
are bit-identical to eager (max abs diff 0.0), as expected for a parameter
the kernel never reads.
Co-authored-by: Copilot <[email protected]>
vaibhavjindal
marked this pull request as draft
August 7, 2026 08:48
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.
Summary
_triton_ropedeclaresbs: tl.constexpr, butbsis never referenced in the kernel body — it appears only in comments. An AST check of the kernel's parameters:Marking a dead parameter
constexprstill forces Triton/Dynamo to specialize on it. That breakstorch.compileas soon as the batch dimension becomes dynamic: the value arrives as aSymInt, and Dynamo trips an internal assertion while tracingLigerRopeFunctionthrough theautograd_function_applyHOP:The failure only surfaces on the second distinct input shape — the first call compiles statically, the second triggers automatic dynamic shapes. That timing made it look like an upstream PyTorch bug rather than a kernel signature issue.
Dropping the
tl.constexprannotation is sufficient. The parameter itself is kept so the call signature and theq size: (bsz, ...)documentation stay intact.Details
Minimal repro (no
transformersdependency), againstmain:Varying
seq_lenalone was always fine, becauseslis an ordinary runtime arg. Only the constexpr batch size was fatal:dynamic=Falsedynamic=Falsedynamic=TrueEnd to end, Qwen3-0.6B with
apply_liger_kernel_to_qwen3(rope=True, ...)now compiles under a plaintorch.compile(model)across all of1x512, 1x2048, 4x512, 4x2048, 8x1024, 8x2048. Previously that required adynamic=Falseworkaround. Explicittorch._dynamo.mark_dynamicon the batch and sequence dims also works now (verified at1x64, 2x128, 8x256).dynamic=Trueremains unsupported, and I don't think it should block this. That flag additionally markshead_dimand the head counts dynamic, which cannot work while block sizing is derived fromtriton.next_power_of_2(head_dim)at launch time. Those dims are fixed by model architecture in practice, so it isn't the case that matters — whereas batch and sequence length genuinely vary, and both work now.Two related observations, deliberately not changed here:
BLOCK_SIZEis also unused in_triton_rope's body. It's harmless today (it's derived from head counts, so it never becomes aSymInt), so I left it rather than widen the diff.qwen2vl_mrope.pyhas the samebs: tl.constexprannotation, but therebsis genuinely used for pointer arithmetic, so it needs its own treatment rather than the same one-line change. The Ascend/NPU backends carry the same pattern and I have no way to test them here.Testing Done
pytest test/transformers/test_rope.py-> 36 passedpytest test/transformers/test_rope.py test/transformers/test_monkey_patch.py-> 98 passedpytest test/convergence/bf16/test_mini_models.py -k "qwen3 or llama or mistral"-> 11 passedmake checkstyle-> passedCompiled vs eager output and gradients are bit-identical (max abs diff
0.0) acrossbs=1,2,4,8— expected, since the kernel never reads this parameter.Hardware Type: H100 80GB HBM3
run
make testto ensure correctnessrun
make checkstyleto ensure code stylerun
make test-convergenceto ensure convergence