Repository navigation
Conversation
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info
📝 Walkthrough
Merge Risk: ⚪ Minimal · up to
|
Co-authored-by: Cursor <cursoragent@cursor.com>
16365ce to
af07316
Compare
|
Updated in
|
| triton = _triton | ||
| tl = _tl | ||
| libdevice = _libdevice | ||
| _MERGED_SWIGLU_FORWARD_KERNEL = _triton.jit(_merged_swiglu_forward_kernel) |
There was a problem hiding this comment.
[P2] On the first compiled training call, this initializer writes module globals and constructs a JITFunction inside the traced custom autograd forward, so fullgraph=True fails unless an eager call has already warmed the cache. Initialize the kernel outside tracing while preserving optional imports, and add a fresh-process compiled-first regression.
There was a problem hiding this comment.
Fixed the compiled-first initialization while keeping the shared helper and fallback for both EP and non-EP consumers. Only the Triton JIT definition is now behind a lazy private-module import. Added fresh-process Inductor regressions without eager warmup; these and the existing MoE and Inductor VJP tests pass.
There was a problem hiding this comment.
[P2] Verified by static inspection: the lazy JIT factory and module-global writes are removed, and _triton_swiglu.py defines the kernel with @triton.jit during import. This addresses the reported cold-start tracing failure; the new subprocess regressions exercise compiled-first calls without eager warmup.
Co-authored-by: Cursor <cursoragent@cursor.com>
What does this PR do?
Fuse the staged
silu(gate) * upactivation in the CUDA merged-FC1 Triton MoE forward into one pointwise kernel. It reads the merged FC1 views directly and writes only the final activation, avoiding the intermediate SiLU tensor and one kernel launch.The SiLU result is rounded to the activation dtype before multiplication. Routing weights remain applied after FC2; saved tensors and backward formulas are unchanged. ROCm and MLU retain the ATen activation path.
Checklist Before Starting
[{modules}] {type}: {description}.Test
Baseline/candidate checks use
38898e9without/with this patch. The fused-model sync selection also passes at the finald02a95ebase.Extended the existing CI-selected
tests/ops/moe_experts/test_moe_experts.pywith exact staged-rounding, input-preservation, empty/strided/clamped-input, and merged-consumer forward/backward comparisons.d02a95ebase.torch.compile(..., backend="eager", fullgraph=True)forward/backward, and CUDA graph forward/backward with replay.make qualityand configured pre-commit checks on the proposed files pass.Training comparison uses the Qwen3 MoE configuration (48 layers, hidden size 2048, expert intermediate size 768, 128 experts, top-k 8), random initialization and BF16. Each dataset uses a nominal sequence budget of 81,920. Baseline, candidate and baseline repeat each run ten updates; the first six are warmup and the last four are measured. Input files, initialization and settings are identical, and effective input lengths match at every compared update. Training measurements use official
38898e9with and without this patch.All ratios below are relative to the first baseline within the same dataset. Peak allocated memory is the candidate peak over the four measured updates, normalized to the corresponding baseline peak.
Candidate throughput improves by about 0.39% on synthetic inputs and 0.98% on documents versus the reported baseline. Every matched steady update is faster, with zero allocation retries in all measured runs. Peak allocated memory is effectively unchanged.
API and Usage Example
No public API, registry selection or configuration changes. The existing
fused_tritonmerged implementation uses the internal helper on its supported path.Design & Code Changes
Checklist Before Submitting
Summary by CodeRabbit
New Features
Tests
Documentation