Add fused RoPE kernel using CuTe DSL

[Copilot]

Implements a fused Rotary Position Embedding (RoPE) operation using NVIDIA's CuTe DSL (CUTLASS), fusing Q and K rotation into a single kernel launch.

Kernel (rope_sm90.cu)

• One CTA per (token, head) pair; threads stride over head_dim/2
• cute::make_tensor / make_gmem_ptr / make_layout for tensor views
• Compile-time HEAD_DIM specialisation via cute::Int<N> (64 / 128 / 256)
• Supports fp16 and bf16; accumulates in float32
• rotate_half convention: x_out = x * cos + rotate_half(x) * sin

Python interface (rope.py)

• Falls back to a pure-PyTorch reference when the CUDA extension is not compiled

Build

• rope_sm90.cu added to the CUDA extension sources in setup.py
• fused_rope_forward registered in bindings.cpp and exported from rynn_scale.ops

Usage

from rynn_scale.ops import fused_rope

# q, k: [total_tokens, num_{q,k}_heads, head_dim]  fp16/bf16
# cos, sin: [total_tokens, head_dim/2]              float32
q_out, k_out = fused_rope(q, k, cos, sin)

Original prompt

帮我用cutedsl写一个fused rope，要求：

创建自 VS Code。

---

🔒 GitHub Advanced Security automatically protects Copilot coding agent pull requests. You can protect all pull requests by enabling Advanced Security for your repositories. Learn more about Advanced Security.

[CLAassistant]

Thank you for your submission! We really appreciate it. Like many open source projects, we ask that you sign our Contributor License Agreement before we can accept your contribution.
_{You have signed the CLA already but the status is still pending? Let us recheck it.}