Add fused DeepSeek-V4 partial interleaved RoPE

[vedavyas2727]

Summary

Add a fused Triton implementation of DeepSeek-V4's partial interleaved rotary position embedding and enable it by default in apply_liger_kernel_to_deepseek_v4.

The kernel matches the Transformers 5.14.1 contract: it rotates only the trailing RoPE slice in FP32, preserves the leading non-rotary channels, accepts the supported unsqueeze_dim layouts and arbitrary positive strides, returns a fresh contiguous tensor, and uses the conjugate transform for backward without saving the input tensor. This also adds the functional/public APIs, monkey-patch and revert coverage, focused correctness tests, a reproducible benchmark script, benchmark data, and the README support-table entry.

Related context: #623. This PR is scoped to DeepSeek-V4 and does not claim to close the broader DeepSeek-V3 request.

Details

DeepSeek-V4 uses adjacent interleaved rotary pairs in only the trailing portion of each head, so the existing Liger Llama-style RoPE kernel is not compatible. The new multi-row kernel handles the production transposed/non-contiguous [B, H, S, D] layout directly and writes contiguous output. It is used for query, shared-KV, and inverse-rotated attention-output calls through the model module's apply_rotary_pos_emb global.

No nn.Module wrapper is added because the upstream surface is a functional model helper. Gradients for cos and sin are intentionally not produced, matching their use as generated positional buffers in DeepSeek-V4.

Testing Done

• Hardware Type: NVIDIA GeForce RTX 5090 Laptop GPU (24 GB, compute capability 12.0), driver 572.76
• Software: Windows, Python 3.13.14, PyTorch 2.11.0+cu128, Triton-Windows 3.7.1, Transformers 5.14.1
• run make test to ensure correctness (targeted operator, patch, and convergence tests run instead)
• run Ruff check and format validation on all touched Python files
• run targeted DeepSeek-V4 fp32 and bf16 convergence, with and without fused linear cross entropy, while asserting the Liger RoPE patch is active

Correctness coverage includes fp16/bf16/fp32/fp64 forward and backward, contiguous and production non-contiguous layouts, unsqueeze_dim=1/2, broadcast and per-batch trigonometric tensors, strided trigonometric tensors, negative sine, odd shapes, full and partial rotary dimensions, functional exports, patch/revert behavior, and instance patching.

Commands/results:

• pytest test/transformers/test_deepseek_v4_rope.py test/transformers/test_monkey_patch.py -k deepseek_v4 -q: 24 passed
• pytest test/transformers/test_monkey_patch.py -q: 62 passed
• pytest test/convergence/fp32/test_mini_models.py -k mini_deepseek_v4 -xvs: 1 passed
• pytest test/convergence/fp32/test_mini_models_with_logits.py -k mini_deepseek_v4 -xvs: 1 passed
• pytest test/convergence/bf16/test_mini_models.py -k mini_deepseek_v4 -xvs: 1 passed
• pytest test/convergence/bf16/test_mini_models_with_logits.py -k mini_deepseek_v4 -xvs: 1 passed

All four convergence suites run with rope=True for DeepSeek-V4 and assert modeling_deepseek_v4.apply_rotary_pos_emb is liger_deepseek_v4_rotary_pos_emb during the Liger path.

Benchmark methodology

B=1, H=64, D=512, rotary dimension R=64, bf16, exact transposed/non-contiguous [B,H,S,D] inputs. Speed uses triton.testing.do_bench with repository quantiles [0.5, 0.2, 0.8]; single calls use 25 ms warmup / 100 ms measurement and the aggregate workload uses 10 ms / 40 ms. Memory is the repository-standard absolute CUDA allocator peak over five runs. Providers are Transformers eager, torch.compile(fullgraph=True), and Liger. Raw results are committed to benchmark/data/all_benchmark_data.csv.

Single-call full forward+backward median latency:

T	HF eager (ms)	torch.compile (ms)	Liger (ms)
512	0.795	0.310	0.329
2048	3.295	1.385	1.458
4096	7.087	2.828	2.687

At T=4096, forward / backward / full medians are:

Provider	Forward (ms)	Backward (ms)	Full (ms)	Full peak (MB)
HF eager	3.229	3.925	7.087	1281
torch.compile	0.904	1.840	2.828	1313
Liger	1.188	1.495	2.687	1025

Aggregate DeepSeek-V4-style Q (H=64) + shared KV (H=1) + inverse output (H=64) full forward+backward:

T	HF eager (ms)	torch.compile (ms)	Liger (ms)
512	2.310	0.856	0.883
2048	7.028	3.498	3.122
4096	15.610	7.023	6.391

At T=4096, aggregate full peak memory is 2321 / 2353 / 2065 MB for HF eager / torch.compile / Liger respectively.

These are operator and three-call attention-block workload measurements, not an end-to-end DeepSeek-V4 training benchmark; this PR does not claim an end-to-end model speedup.