[GDN] Split gdn attention into conv1d and gated_delta_rule#402
Open
HeJunyan wants to merge 9 commits into
Open
[GDN] Split gdn attention into conv1d and gated_delta_rule#402HeJunyan wants to merge 9 commits into
HeJunyan wants to merge 9 commits into
Conversation
HeJunyan
changed the title
Contributor
There was a problem hiding this comment.
Pull request overview
Note
Copilot was unable to run its full agentic suite in this review.
This PR splits the legacy fused XPU gdn_attention op into two explicit ops (causal_conv1d and gated_delta_rule), updates test coverage to validate the split-path behavior, and adds standalone benchmarks for each stage.
Changes:
- Added new Torch ops/bindings for
causal_conv1dandgated_delta_ruleand refactored the legacygdn_attentionto be a wrapper over them. - Updated existing GDN attention tests to call the split ops directly and added new dedicated tests for each stage.
- Updated and added benchmarks to measure fused vs. split-stage performance.
Reviewed changes
Copilot reviewed 10 out of 10 changed files in this pull request and generated 7 comments.
Show a summary per file
| File | Description |
|---|---|
csrc/xpu/torch_bindings.cpp |
Registers new Torch op schemas and implementations for the two split ops. |
csrc/xpu/ops.h |
Exposes new C++ op APIs for causal_conv1d and gated_delta_rule. |
csrc/xpu/gdn_attn/gdn_attn_interface.cpp |
Implements the split ops and rewrites legacy gdn_attention as a wrapper. |
tests/gdn_attn/test_gdn_attn_padded.py |
Updates padded-leading-dim test to call split ops instead of fused op. |
tests/gdn_attn/test_gdn_attn.py |
Adds split-op validation path and adds a “legacy fused op” backward-compat test. |
tests/gdn_attn/test_causal_conv1d.py |
New test file validating conv-stage outputs against fused reference. |
tests/gdn_attn/test_gated_delta_rule.py |
New test file validating delta-stage outputs against fused reference. |
benchmark/benchmark_gdn_attn.py |
Updates fused benchmark runner logic to execute split ops. |
benchmark/benchmark_causal_conv1d.py |
New benchmark for conv stage (with bandwidth/MBU/TFLOPS reporting). |
benchmark/benchmark_gated_delta_rule.py |
New benchmark for delta stage (with bandwidth/MBU/TFLOPS reporting). |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Wei-Lin-Intel
suggested changes
Jun 9, 2026
HeJunyan
force-pushed
the
split_gdn_attention
branch
4 times, most recently
from
June 10, 2026 10:47
8a8f7f9 to
dd30d2a
Compare
Replace the fused torch.ops._xpu_C.gdn_attention op with two exported ops:
- causal_conv1d: runs the conv stage, writes z, updates conv_state, and
returns the intermediate {q,k,v,b,a} tensors (spec buffers first, then
non-spec buffers; native vs XE2 prefill shapes preserved).
- gated_delta_rule: consumes those intermediates, recomputes the same path
decisions, advances ssm_state, and writes core_attn_out.
Both ops preserve the active-prefix narrowing for cudagraph-padded inputs.
Signed-off-by: He Junyan <junyan.he@intel.com>
…delta_rule ops
Replace the single torch.ops._xpu_C.gdn_attention call in test_gdn_attn.py
(non-spec and spec/mtp paths), test_gdn_attn_padded.py (padded + reference
calls), and benchmark_gdn_attn.py with the two-op sequence: causal_conv1d
returns the intermediate {q,k,v,b,a} list which is forwarded to
gated_delta_rule. Behavior and assertions are unchanged.
Signed-off-by: He Junyan <junyan.he@intel.com>
Add test_causal_conv1d.py and test_gated_delta_rule.py mirroring test_gdn_attn.py (same shapes, params, and golden references): - test_causal_conv1d[/_mtp]: drives only the causal_conv1d op and asserts the conv-stage outputs (z and conv_state) against the fused reference. - test_gated_delta_rule[/_mtp]: feeds causal_conv1d intermediates into gated_delta_rule and asserts the delta-stage outputs (core_attn_out and ssm_state) against the fused reference. Both files reuse ref_gdn_attention / ref_gdn_attention_spec from test_gdn_attn.py to avoid duplicating the reference implementation. Signed-off-by: He Junyan <junyan.he@intel.com>
Add benchmark_causal_conv1d.py and benchmark_gated_delta_rule.py, imitating benchmark_gdn_attn.py (same triton.testing.perf_report harness, model shapes, and workloads, reused via import to stay in lockstep): - benchmark_causal_conv1d times only the conv stage, with byte/FLOP models covering the conv1d read/write set (projected states, conv_state, z and the q/k/v/b/a intermediates). - benchmark_gated_delta_rule runs the conv stage once outside the timing loop to produce the intermediates, then times only gated_delta_rule, with byte/FLOP models covering the SSM recurrence read/write set. Both report latency (us), memory bandwidth (GB/s), MBU (%) and TFLOPS. Signed-off-by: He Junyan <junyan.he@intel.com>
Re-add gdn_attention as a thin wrapper that chains causal_conv1d then gated_delta_rule, preserving the original fused API. Add test_gdn_attention_legacy as a full copy of the original test_gdn_attention case (same parametrization), driving the fused op to verify backward-compat coverage matches the split-op test. Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Signed-off-by: He Junyan <junyan.he@intel.com>
Remove the opaque intermediates list parameter from gated_delta_rule and replace it with an explicit single group of q,k,v,b,a tensors. causal_conv1d still returns the vector of intermediates; callers now unpack it (*intermediates) into the explicit params. The legacy gdn_attention wrapper unpacks the vector into the 5 args. Updated ops.h, torch_bindings.cpp schema, all Python tests and benchmarks accordingly. Signed-off-by: He Junyan <junyan.he@intel.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
According to comments: 1. Add more TORCH_CHECK for operators. 2. Improve the test_causal_conv1d and test_gated_delta_rule, using standalone test reference functions. Signed-off-by: He Junyan <junyan.he@intel.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
HeJunyan
force-pushed
the
split_gdn_attention
branch
from
June 30, 2026 10:20
dd30d2a to
4293708
Compare
The XE2 prefill conv path fuses the downstream l2norm into q/k (and scales q by 1/sqrt(head_k_dim)), while the decode path emits raw conv output and defers l2norm to gated_delta_rule. The dedicated conv test's prefill-layout assertion compared raw ref_q/ref_k against the already-normalized kernel output, causing all prefill/mix_mode cases to fail. Apply the same l2norm and scale to the reference before comparison (mirroring ref_gdn_attention). Signed-off-by: He Junyan <junyan.he@intel.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
HeJunyan
force-pushed
the
split_gdn_attention
branch
from
July 1, 2026 07:28
8a79c6e to
7591532
Compare
Signed-off-by: He Junyan <junyan.he@intel.com>
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.
The current gdn_attention includes everything and is too huge. It is inconvenient for profiling. So we split it into conv1d and gated_delta_rule. We also add test cases and benchmark for conv1d and gated_delta_rule separately. For the legacy usage, we still keep the original gdn_attention operator, which is a simple wrap of calling conv1d and gated_delta_rule.