Skip to content

[RWKV7] Drop some kernels to enhance speed#346

Merged
zhiyuan1i merged 10 commits intomainfrom
rwkv-speedup
Apr 11, 2025
Merged

[RWKV7] Drop some kernels to enhance speed#346
zhiyuan1i merged 10 commits intomainfrom
rwkv-speedup

Conversation

@zhiyuan1i
Copy link
Copy Markdown
Collaborator

@zhiyuan1i zhiyuan1i commented Apr 11, 2025

Summary by CodeRabbit

  • Refactor

    • Streamlined several internal routines to optimize tensor operations, resulting in simplified control flows and updated default parameters for improved precision handling and chunking behavior.
    • Enhanced error messaging to better guide users on supported data types.
  • Tests

    • Updated test configurations to leverage enhanced precision settings.

@coderabbitai
Copy link
Copy Markdown
Contributor

coderabbitai bot commented Apr 11, 2025

Caution

Review failed

The pull request is closed.

Walkthrough

This pull request refactors several functions across multiple modules by removing the input_precision parameter from RWKV7-related functions and updating their signatures. It also adds a warning in the delta rule function for unsupported float32 inputs, adjusts kernel loops and indexing logic, and reduces several default chunk_size values from 64 to 16. Other changes include streamlining tensor masking operations, introducing a new DTYPE parameter for flexible data type handling in fused addcmul kernels, and updating test parameters accordingly.

Changes

File(s) Change Summary
fla/layers/rwkv7.py
fla/ops/rwkv7/chunk.py
fla/ops/rwkv7/fused_recurrent.py
Removed input_precision parameter and updated function signatures in RWKV7 functions.
fla/ops/generalized_delta_rule/dplr/chunk.py Added a warning in chunk_dplr_delta_rule to alert users when using torch.float32, advising the use of bfloat16.
fla/ops/generalized_delta_rule/dplr/chunk_A_bwd.py
fla/ops/generalized_delta_rule/dplr/chunk_A_fwd.py
Simplified backward and forward kernel logic by removing the NC parameter, adjusting indexing, and streamlining loop conditions.
fla/ops/generalized_delta_rule/dplr/chunk_h_bwd.py
fla/ops/generalized_delta_rule/dplr/chunk_o_bwd.py
fla/ops/generalized_delta_rule/dplr/chunk_o_fwd.py
Updated default chunk_size from 64 to 16 and revised the masking operations for tensor processing.
fla/ops/generalized_delta_rule/dplr/fused_recurrent.py Removed an unnecessary blank line following a warning message.
fla/ops/generalized_delta_rule/dplr/wy_fast_bwd.py Introduced a precomputed boolean mask variable (m_i) to improve the clarity of masking operations.
fla/ops/rwkv7/fused_addcmul.py Added a new DTYPE parameter for flexible data type handling and an inplace parameter to the addcmul_bwd1 function.
tests/ops/test_rwkv7.py Updated the test parameterization by replacing torch.float16 with torch.bfloat16.

Sequence Diagram(s)

sequenceDiagram
    participant Test as Test Suite
    participant FRwkv7 as fused_recurrent_rwkv7
    participant RWkv7 as rwkv7_fn
    Test->>FRwkv7: Initiate forward pass
    FRwkv7->>RWkv7: Call without input_precision
    RWkv7-->>FRwkv7: Return computed tensor
    FRwkv7-->>Test: Pass along output
Loading
sequenceDiagram
    participant Caller
    participant Func as chunk_dplr_delta_rule
    Caller->>Func: Invoke function with tensor q (float32)
    Func->>Func: Check tensor data type
    alt Tensor is float32
       Func->>Caller: Emit warning ("ChunkDeltaRuleFunction does not support float32...")
    else
       Func->>Caller: Proceed with computation
    end
Loading

Possibly related PRs

Suggested reviewers

  • yzhangcs

Poem

I’m Bunny, hopping through refactored code with delight,
Input precision removed, now everything’s light.
Warnings and kernels now dance in a cleaner phase,
Chunk sizes shrink and masking sings in perfect ways.
In the burrow of code, every change makes me cheer—
Hoppy coding adventures lie delightfully near!
🐇💻

Tip

⚡💬 Agentic Chat (Pro Plan, General Availability)
  • We're introducing multi-step agentic chat in review comments and issue comments, within and outside of PR's. This feature enhances review and issue discussions with the CodeRabbit agentic chat by enabling advanced interactions, including the ability to create pull requests directly from comments and add commits to existing pull requests.

📜 Recent review details

Configuration used: CodeRabbit UI
Review profile: CHILL
Plan: Pro

📥 Commits

Reviewing files that changed from the base of the PR and between 2f79d88 and 11ac5bd.

📒 Files selected for processing (13)
  • fla/layers/rwkv7.py (0 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk.py (1 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk_A_bwd.py (9 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk_A_fwd.py (5 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk_h_bwd.py (1 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk_o_bwd.py (6 hunks)
  • fla/ops/generalized_delta_rule/dplr/chunk_o_fwd.py (2 hunks)
  • fla/ops/generalized_delta_rule/dplr/fused_recurrent.py (0 hunks)
  • fla/ops/generalized_delta_rule/dplr/wy_fast_bwd.py (2 hunks)
  • fla/ops/rwkv7/chunk.py (0 hunks)
  • fla/ops/rwkv7/fused_addcmul.py (8 hunks)
  • fla/ops/rwkv7/fused_recurrent.py (0 hunks)
  • tests/ops/test_rwkv7.py (1 hunks)
✨ Finishing Touches
  • 📝 Generate Docstrings

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share
🪧 Tips

Chat

There are 3 ways to chat with CodeRabbit:

  • Review comments: Directly reply to a review comment made by CodeRabbit. Example:
    • I pushed a fix in commit <commit_id>, please review it.
    • Generate unit testing code for this file.
    • Open a follow-up GitHub issue for this discussion.
  • Files and specific lines of code (under the "Files changed" tab): Tag @coderabbitai in a new review comment at the desired location with your query. Examples:
    • @coderabbitai generate unit testing code for this file.
    • @coderabbitai modularize this function.
  • PR comments: Tag @coderabbitai in a new PR comment to ask questions about the PR branch. For the best results, please provide a very specific query, as very limited context is provided in this mode. Examples:
    • @coderabbitai gather interesting stats about this repository and render them as a table. Additionally, render a pie chart showing the language distribution in the codebase.
    • @coderabbitai read src/utils.ts and generate unit testing code.
    • @coderabbitai read the files in the src/scheduler package and generate a class diagram using mermaid and a README in the markdown format.
    • @coderabbitai help me debug CodeRabbit configuration file.

Note: Be mindful of the bot's finite context window. It's strongly recommended to break down tasks such as reading entire modules into smaller chunks. For a focused discussion, use review comments to chat about specific files and their changes, instead of using the PR comments.

CodeRabbit Commands (Invoked using PR comments)

  • @coderabbitai pause to pause the reviews on a PR.
  • @coderabbitai resume to resume the paused reviews.
  • @coderabbitai review to trigger an incremental review. This is useful when automatic reviews are disabled for the repository.
  • @coderabbitai full review to do a full review from scratch and review all the files again.
  • @coderabbitai summary to regenerate the summary of the PR.
  • @coderabbitai generate docstrings to generate docstrings for this PR.
  • @coderabbitai resolve resolve all the CodeRabbit review comments.
  • @coderabbitai plan to trigger planning for file edits and PR creation.
  • @coderabbitai configuration to show the current CodeRabbit configuration for the repository.
  • @coderabbitai help to get help.

Other keywords and placeholders

  • Add @coderabbitai ignore anywhere in the PR description to prevent this PR from being reviewed.
  • Add @coderabbitai summary to generate the high-level summary at a specific location in the PR description.
  • Add @coderabbitai anywhere in the PR title to generate the title automatically.

CodeRabbit Configuration File (.coderabbit.yaml)

  • You can programmatically configure CodeRabbit by adding a .coderabbit.yaml file to the root of your repository.
  • Please see the configuration documentation for more information.
  • If your editor has YAML language server enabled, you can add the path at the top of this file to enable auto-completion and validation: # yaml-language-server: $schema=https://coderabbit.ai/integrations/schema.v2.json

Documentation and Community

  • Visit our Documentation for detailed information on how to use CodeRabbit.
  • Join our Discord Community to get help, request features, and share feedback.
  • Follow us on X/Twitter for updates and announcements.

yzhangcs and others added 6 commits April 11, 2025 08:01
* [Attn] Remove `head_first` & rename `offsets` to `cu_seqlens`

* Delete 256 headdim tests

* [DeltaNet] Remove `head_first` & rename `offsets` to `cu_seqlens`

* Rename `offsets` to `cu_seqlens` across GLA/GSA/RWKV6

* Fix NSA tests

* Fix TTT checks

* Rename `offsets` to `cu_seqlens` in DPLR and IPLR implementations, updating related logic and heuristics for variable-length sequences.
@zhiyuan1i zhiyuan1i marked this pull request as ready for review April 11, 2025 08:21
@zhiyuan1i zhiyuan1i requested a review from Copilot April 11, 2025 08:21
Copy link
Copy Markdown

Copilot AI left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot reviewed 13 out of 13 changed files in this pull request and generated no comments.

Comments suppressed due to low confidence (4)

tests/ops/test_rwkv7.py:18

  • Consider adding tests for torch.float16 if needed to cover alternative precision paths, or ensure that bfloat16 is the exclusively supported type.
@pytest.mark.parametrize("dtype", [torch.bfloat16])

fla/ops/rwkv7/fused_recurrent.py:22

  • Ensure that removing the input_precision parameter does not break downstream calls expecting this parameter; update documentation if necessary.
input_precision: Optional[torch.dtype] = torch.bfloat16,

fla/ops/rwkv7/fused_addcmul.py:55

  • Confirm that replacing the explicit conversion to tl.float32 with conversion to DTYPE maintains numerical stability and performance; verify that the DTYPE parameter is set appropriately at all call sites.
b_hiddn = tl.load(hidden_ptr + (xindex), xmask, other=0.).to(DTYPE)

fla/ops/generalized_delta_rule/dplr/chunk.py:333

  • [nitpick] The warning message for float32 support could be enhanced by providing explicit guidance on alternative behavior or raising an error for unsupported dtypes.
if q.dtype == torch.float32:

@zhiyuan1i zhiyuan1i merged commit a1b163c into main Apr 11, 2025
3 of 6 checks passed
@zhiyuan1i zhiyuan1i deleted the rwkv-speedup branch April 11, 2025 08:22
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants