Publish RWKV7-1.5B-20260805

.gitattributes

@@ -1,35 +1,3 @@
-*.7z filter=lfs diff=lfs merge=lfs -text
-*.arrow filter=lfs diff=lfs merge=lfs -text
-*.bin filter=lfs diff=lfs merge=lfs -text
-*.bz2 filter=lfs diff=lfs merge=lfs -text
-*.ckpt filter=lfs diff=lfs merge=lfs -text
-*.ftz filter=lfs diff=lfs merge=lfs -text
-*.gz filter=lfs diff=lfs merge=lfs -text
-*.h5 filter=lfs diff=lfs merge=lfs -text
-*.joblib filter=lfs diff=lfs merge=lfs -text
-*.lfs.* filter=lfs diff=lfs merge=lfs -text
-*.mlmodel filter=lfs diff=lfs merge=lfs -text
-*.model filter=lfs diff=lfs merge=lfs -text
-*.msgpack filter=lfs diff=lfs merge=lfs -text
-*.npy filter=lfs diff=lfs merge=lfs -text
-*.npz filter=lfs diff=lfs merge=lfs -text
-*.onnx filter=lfs diff=lfs merge=lfs -text
-*.ot filter=lfs diff=lfs merge=lfs -text
-*.parquet filter=lfs diff=lfs merge=lfs -text
-*.pb filter=lfs diff=lfs merge=lfs -text
-*.pickle filter=lfs diff=lfs merge=lfs -text
-*.pkl filter=lfs diff=lfs merge=lfs -text
-*.pt filter=lfs diff=lfs merge=lfs -text
-*.pth filter=lfs diff=lfs merge=lfs -text
-*.rar filter=lfs diff=lfs merge=lfs -text
 *.safetensors filter=lfs diff=lfs merge=lfs -text
-saved_model/**/* filter=lfs diff=lfs merge=lfs -text
-*.tar.* filter=lfs diff=lfs merge=lfs -text
-*.tar filter=lfs diff=lfs merge=lfs -text
-*.tflite filter=lfs diff=lfs merge=lfs -text
-*.tgz filter=lfs diff=lfs merge=lfs -text
-*.wasm filter=lfs diff=lfs merge=lfs -text
-*.xz filter=lfs diff=lfs merge=lfs -text
-*.zip filter=lfs diff=lfs merge=lfs -text
-*.zst filter=lfs diff=lfs merge=lfs -text
-*tfevents* filter=lfs diff=lfs merge=lfs -text
+*.pth filter=lfs diff=lfs merge=lfs -text
+*.bin filter=lfs diff=lfs merge=lfs -text

LICENSE

@@ -0,0 +1,201 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and do not
+          modify the License. You may add Your own attribution notices
+          within Derivative Works that You distribute, alongside or as an
+          addendum to the NOTICE text from the Work, provided that such
+          additional attribution notices cannot be construed as modifying
+          the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer, and
+      charge a fee for, acceptance of support, warranty, indemnity, or
+      other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

NOTICE

@@ -0,0 +1,4 @@
+RWKV-7 model release
+Source: BlinkDL/rwkv7-g1/rwkv7-g1i-1.5b-20260805-ctx16384.pth
+Bundled Transformers RWKV-7 code: huggingface/transformers@4ad9ed0747ed6ba75c787e8f9040dcd64b166ee2 (Apache-2.0).
+Exported inference bundle licensed under Apache-2.0.

README.md

@@ -0,0 +1,279 @@
+---
+library_name: transformers
+pipeline_tag: text-generation
+license: apache-2.0
+language:
+  - "en"
+  - "zh"
+  - "fr"
+  - "es"
+  - "de"
+  - "pt"
+  - "ru"
+  - "it"
+  - "ja"
+  - "ko"
+  - "vi"
+  - "ar"
+datasets:
+  - "HuggingFaceFW/fineweb-edu"
+  - "mlfoundations/dclm-baseline-1.0"
+  - "cerebras/SlimPajama-627B"
+  - "EleutherAI/pile"
+  - "bigcode/starcoderdata"
+  - "oscar-corpus/OSCAR-2301"
+tags:
+  - rwkv
+  - rwkv7
+  - recurrent
+  - causal-lm
+  - conversational
+---
+
+<!-- markdownlint-disable first-line-h1 -->
+<!-- markdownlint-disable html -->
+
+<div align="center">
+  <a href="https://www.rwkv.com/">
+    <img src="https://www.rwkv.com/images/avatar.png" width="140" alt="RWKV logo" />
+  </a>
+  <h1>RWKV7-1.5B-20260805</h1>
+  <p><strong>RWKV-7 “Goose” · constant-state recurrent language modeling</strong></p>
+</div>
+
+<div align="center">
+  <a href="https://www.rwkv.com/"><img alt="Website" src="https://img.shields.io/badge/Website-RWKV-16a7c9" /></a>
+  <a href="https://huggingface.co/BlinkDL"><img alt="Hugging Face" src="https://img.shields.io/badge/%F0%9F%A4%97%20Hugging%20Face-BlinkDL-ffc107" /></a>
+  <a href="https://github.com/BlinkDL/RWKV-LM"><img alt="GitHub" src="https://img.shields.io/badge/GitHub-RWKV--LM-181717?logo=github" /></a>
+  <a href="https://arxiv.org/abs/2503.14456v2"><img alt="RWKV-7 paper" src="https://img.shields.io/badge/Paper-arXiv%3A2503.14456-b31b1b" /></a>
+  <a href="LICENSE"><img alt="License" src="https://img.shields.io/badge/License-apache-2.0-4c8bf5" /></a>
+</div>
+
+---
+
+## Model introduction
+
+This is an official BlinkDL release of **RWKV-7 Goose** in Hugging Face
+Transformers format. RWKV-7 is an attention-free recurrent architecture with a
+constant-size recurrent state and constant inference work per generated token.
+Training remains parallelizable.
+
+This checkpoint is a **base model** pretrained with web, code, synthetic, instruction, chat, and reasoning data. It is suitable for evaluation, post-training, and fine-tuning; the included chat template is a prompt interface, not a claim that the checkpoint is a safety-aligned assistant.
+
+The Transformers integration, conversion, release packaging, Fast Tokenizer, and
+optional TileLang inference implementation are distributed with this release.
+
+## Highlights
+
+- **Constant recurrent state:** memory does not grow like an attention KV cache.
+- **Bundled Transformers integration:** auditable remote configuration and modeling
+  modules provide generation, recurrent cache continuation, training, and LoRA
+  workflows on Transformers 5.15+.
+- **Exact Fast Tokenizer:** self-contained Rust-backed `tokenizer.json`, generated
+  from the canonical RWKV World byte vocabulary during conversion.
+- **Chat-ready:** `chat_template.jinja` supports system, multi-turn, thinking, and
+  strict model-generated tool-call prompts.
+- **Optional optimized runtime:** the isolated [`inference/`](inference/) bundle
+  provides PyTorch fallback and TileLang acceleration without changing the
+  standard model root.
+
+## Model overview
+
+| Field | Value |
+| --- | --- |
+| Repository | `aabbdev/RWKV7-1.5B-20260805` |
+| Architecture class | `Rwkv7ForCausalLM` |
+| Public size label | `1.5`B |
+| Source parameters | `1,527,668,736` |
+| Serialized parameters | `1,527,668,736` |
+| Synthesized compatibility tensors | `0` |
+| Layers | `24` |
+| Hidden / FFN size | `2048` / `8192` |
+| Heads / head size | `32` / `64` |
+| Vocabulary | `65536` |
+| Training context | `16384 tokens` |
+| Weight dtype | `bfloat16` |
+| Numerical conversion | `source dtype preserved` |
+| Metadata profile | `g1i` |
+| Metadata provenance | `locked-profile` |
+| Source checkpoint | [`BlinkDL/rwkv7-g1/rwkv7-g1i-1.5b-20260805-ctx16384.pth`](https://huggingface.co/BlinkDL/rwkv7-g1/blob/ede85bf8ab2e59aff7d7ca909fbbc73317866d89/rwkv7-g1i-1.5b-20260805-ctx16384.pth) |
+| Source SHA-256 | `32ef7b5bf4dc8bde843cf26dfad809a1f527e2e76a9e790e7d406e71bcd785da` |
+
+## Transformers quickstart
+
+The repository includes `configuration_rwkv7.py` and `modeling_rwkv7.py`, adapted
+from the Transformers RWKV-7 integration at commit
+[`4ad9ed0`](https://github.com/huggingface/transformers/commit/4ad9ed0747ed6ba75c787e8f9040dcd64b166ee2).
+Review those files and pin a model-repository revision in production. Passing
+`trust_remote_code=True` selects this bundled implementation even when the local
+Transformers installation also provides native RWKV-7 support.
+
+```python
+import torch
+from transformers import (
+    AutoModelForCausalLM,
+    AutoTokenizer,
+    PreTrainedConfig,
+)
+
+model_id = "aabbdev/RWKV7-1.5B-20260805"
+tokenizer = AutoTokenizer.from_pretrained(
+    model_id,
+    config=PreTrainedConfig(),
+)
+model = AutoModelForCausalLM.from_pretrained(
+    model_id,
+    trust_remote_code=True,
+    dtype=torch.bfloat16,
+)
+```
+
+The recurrent cache returned by the model can be passed back for incremental
+decoding. Use an `attention_mask` for padded batches.
+
+## Chat quickstart
+
+```python
+import re
+
+import torch
+from transformers import AutoModelForCausalLM, AutoTokenizer, PreTrainedConfig
+
+
+THINK_RE = re.compile(r"\A<think>?\s*(.*?)\s*</think>?", re.DOTALL)
+
+
+def assistant_content(completion, thinking, *, close_incomplete=False):
+    prefix = "<think" if thinking else "<think></think"
+    reply = prefix + completion
+    thinking_block = THINK_RE.match(reply)
+    if thinking:
+        if thinking_block is not None or not close_incomplete:
+            return reply.strip()
+        return f"{reply.rstrip()}\n</think>".strip()
+    return "" if thinking_block is None else reply[thinking_block.end():].strip()
+
+model_id = "aabbdev/RWKV7-1.5B-20260805"
+tokenizer = AutoTokenizer.from_pretrained(
+    model_id,
+    config=PreTrainedConfig(),
+)
+model = AutoModelForCausalLM.from_pretrained(
+    model_id,
+    trust_remote_code=True,
+    dtype=torch.bfloat16,
+).to("cuda")
+
+messages = [{"role": "user", "content": "Explain why RWKV uses constant state."}]
+thinking = False
+max_new_tokens = 256
+inputs = tokenizer.apply_chat_template(
+    messages,
+    tokenize=True,
+    add_generation_prompt=True,
+    thinking=thinking,
+    return_dict=True,
+    return_tensors="pt",
+).to(model.device)
+
+output = model.generate(
+    **inputs,
+    max_new_tokens=max_new_tokens,
+    do_sample=True,
+    temperature=1.0,
+    top_p=0.5,
+    eos_token_id=0,
+    pad_token_id=0,
+    stop_strings=["\n\nUser:"],
+    tokenizer=tokenizer,
+)
+completion = tokenizer.decode(
+    output[0, inputs["input_ids"].shape[1]:],
+    skip_special_tokens=True,
+)
+completion = completion.split("\n\nUser:", 1)[0]
+reached_token_limit = output.shape[1] - inputs["input_ids"].shape[1] >= max_new_tokens
+print(
+    assistant_content(
+        completion,
+        thinking,
+        close_incomplete=reached_token_limit,
+    )
+)
+```
+
+Set `thinking=True` for the RWKV thinking prefix. The intentional generation
+prefixes are `Assistant: <think></think` and `Assistant: <think`; do not append a
+closing `>` to them. The post-processing above reconstructs that prefix before
+removing an empty thinking block or preserving an enabled one. If generation hits
+the token limit inside thinking, it closes the displayed block before returning it.
+Reference stops are token ID `0` and `\n\nUser:`.
+
+Strip trailing spaces from user input. The official RWKV prompt guide is available
+in [`RWKV7-G1x-templates.txt`](https://github.com/BlinkDL/RWKV-LM/blob/main/RWKV-v7/RWKV7-G1x-templates.txt).
+
+## Optimized local inference
+
+Install the versions listed in `inference/requirements.txt`, then run the bundled
+interactive chat:
+
+```bash
+python inference/generate.py --model aabbdev/RWKV7-1.5B-20260805 --backend auto --interactive
+```
+
+Or independent prompts separated by blank lines:
+
+```bash
+python inference/generate.py \
+  --model aabbdev/RWKV7-1.5B-20260805 \
+  --backend auto \
+  --input-file prompts.txt
+```
+
+`--backend auto` uses validated exact optimized boundaries and otherwise falls
+back to PyTorch. Full explicit TileLang execution can change floating-point
+operation order and requires checkpoint-, dtype-, shape-, and device-specific
+parity validation.
+
+## Tokenizer
+
+The model root contains one self-contained tokenizer artifact: `tokenizer.json`.
+Textual `vocab.json` and `rwkv_vocab_v20230424.txt` files are intentionally omitted
+because they would duplicate the tokenizer used by Transformers. The tokenizer is
+loaded natively as `PreTrainedTokenizerFast` and never executes remote Python code.
+The explicit generic config prevents `AutoTokenizer` from probing the remote model
+configuration and emitting a harmless model-type fallback warning.
+
+## Intended use and limitations
+
+- This is a base causal language model. Quality, instruction following, and
+  language behavior depend on the checkpoint and downstream prompting or
+  post-training.
+- Assisted or speculative decoding that requires recurrent-cache rollback is not
+  supported without retaining prior state snapshots.
+- Optimized support depends on GPU architecture, dtype, batch, and shape.
+  Unsupported `auto` configurations fall back to pure PyTorch.
+- Explicit full TileLang execution can change floating-point operation order and
+  requires checkpoint-, dtype-, shape-, and device-specific parity validation.
+- No safety, bias, toxicity, factuality, or high-stakes-use evaluation is claimed
+  by this model card.
+
+## License and provenance
+
+The model weights use the locked profile license `apache-2.0`. The exported inference bundle is licensed separately under [Apache-2.0](LICENSE). The bundled Transformers configuration and modeling modules
+retain their Apache-2.0 headers. See [`NOTICE`](NOTICE) and the source checkpoint
+link above for provenance.
+
+## Citation
+
+```bibtex
+@misc{peng2025250314456,
+  title         = {RWKV-7 "Goose" with Expressive Dynamic State Evolution},
+  author        = {Bo Peng and Ruichong Zhang and Daniel Goldstein and Eric Alcaide and Xingjian Du and Haowen Hou and Jiaju Lin and Jiaxing Liu and Janna Lu and William Merrill and Guangyu Song and Kaifeng Tan and Saiteja Utpala and Nathan Wilce and Johan S. Wind and Tianyi Wu and Daniel Wuttke and Christian Zhou-Zheng},
+  year          = {2025},
+  eprint        = {2503.14456v2},
+  archivePrefix = {arXiv},
+  primaryClass  = {cs.CL},
+  url           = {https://arxiv.org/abs/2503.14456v2},
+}
+```

chat_template.jinja

@@ -0,0 +1,56 @@
+{%- set add_generation_prompt = add_generation_prompt | default(true) -%}
+{%- set thinking = thinking | default(false) -%}
+{%- set bos_token = bos_token | default('', true) -%}
+{%- set tools = tools | default([], true) -%}
+{%- set ns = namespace(system_prompt='') -%}
+{%- for message in messages -%}
+    {%- if message.role == 'system' -%}
+        {%- set ns.system_prompt = message.content | trim -%}
+    {%- endif -%}
+{%- endfor -%}
+{{- bos_token -}}
+{%- if ns.system_prompt or tools | length > 0 -%}
+{{ 'System: ' }}{{ ns.system_prompt }}
+{%- if tools | length > 0 -%}
+{%- if ns.system_prompt %}{{ '\n' }}{%- endif -%}
+{{ 'Tools:\n' -}}
+{{ tools | tojson }}
+{{ '\nWhen using a tool, return only a compact JSON function call in a ```json block, like {"name":"calculator","arguments":{"expression":"2+2"}}. The `name` field must be top-level, never inside `arguments`. Do not copy the tool schema into arguments. Otherwise answer normally.' }}
+{%- endif -%}
+{{ '\n\n' }}
+{%- endif -%}
+{%- for message in messages -%}
+{%- if message.role == 'user' -%}
+{{ 'User: ' ~ (message.content | trim) ~ '\n\n' }}
+{%- elif message.role == 'assistant' -%}
+{%- set content = message.content | default('', true) | trim -%}
+{{ 'Assistant:' }}
+{%- if message.tool_calls is defined and message.tool_calls | length > 0 -%}
+{%- if content %}{{ ' ' ~ content ~ '\n' }}{%- endif -%}
+{%- for tool_call in message.tool_calls -%}
+    {%- if tool_call.function is defined -%}
+        {%- set name = tool_call.function.name | default('') -%}
+        {%- set args = tool_call.function.arguments | default({}, true) -%}
+    {%- else -%}
+        {%- set name = tool_call.name | default('') -%}
+        {%- set args = tool_call.arguments | default({}, true) -%}
+    {%- endif -%}
+{{ ' ```json\n' -}}
+{{ '{"name": ' }}{{ name | tojson }}{{ ', "arguments": ' }}{% if args is string %}{{ args }}{% else %}{{ args | tojson }}{% endif %}{{ '}\n' -}}
+{{ '```' }}{{ '\n' if not loop.last else '' }}
+{%- endfor -%}
+{%- elif content -%}
+{{ ' ' ~ content }}
+{%- endif -%}
+{{ '\n\n' }}
+{%- elif message.role == 'tool' -%}
+{{ 'User: Function output:\n' ~ (message.content | trim) ~ '\n\n' }}
+{%- endif -%}
+{%- endfor -%}
+{%- if add_generation_prompt -%}
+{%- if thinking -%}
+{{ 'Assistant: <think' }}
+{%- else -%}
+{{ 'Assistant: <think></think' }}
+{%- endif -%}
+{%- endif -%}

config.json

@@ -0,0 +1,31 @@
+{
+  "a_low_rank_dim": 96,
+  "architectures": [
+    "Rwkv7ForCausalLM"
+  ],
+  "auto_map": {
+    "AutoConfig": "configuration_rwkv7.Rwkv7Config",
+    "AutoModel": "modeling_rwkv7.Rwkv7Model",
+    "AutoModelForCausalLM": "modeling_rwkv7.Rwkv7ForCausalLM"
+  },
+  "bos_token_id": 0,
+  "decay_low_rank_dim": 96,
+  "dtype": "bfloat16",
+  "eos_token_id": 0,
+  "gate_low_rank_dim": 256,
+  "head_dim": 64,
+  "hidden_size": 2048,
+  "intermediate_size": 8192,
+  "model_type": "rwkv7",
+  "norm_bias": true,
+  "norm_eps": 1e-05,
+  "num_heads": 32,
+  "num_hidden_layers": 24,
+  "pad_token_id": 0,
+  "tie_word_embeddings": false,
+  "use_cache": true,
+  "v_low_rank_dim": 64,
+  "vocab_size": 65536,
+  "wkv_implementation": "eager",
+  "wkv_state_dtype": "float32"
+}

configuration_rwkv7.py

@@ -0,0 +1,153 @@
+# Copyright 2026 The RWKV team and The HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""RWKV-7 (Goose) model configuration."""
+
+from huggingface_hub.dataclasses import strict
+
+from transformers.configuration_utils import PreTrainedConfig
+from transformers.utils import auto_docstring
+
+
+@auto_docstring(
+    checkpoint="RWKV/RWKV7-1.5B-20260805",
+    custom_intro="""
+    Configuration for [`Rwkv7Model`], an all-recurrent (attention-free) RWKV-7 "Goose"
+    model. Instantiating with the defaults yields the ~0.1B RWKV-7 configuration.
+
+    Parameter names follow the upstream RWKV reference implementation
+    (`BlinkDL/RWKV-LM`) rather than a renamed variant.
+    """,
+)
+@strict
+class Rwkv7Config(PreTrainedConfig):
+    r"""
+    vocab_size (`int`, *optional*, defaults to 65536):
+        Vocabulary size (RWKV "world" tokenizer).
+    hidden_size (`int`, *optional*, defaults to 768):
+        Model width `C`.
+    num_hidden_layers (`int`, *optional*, defaults to 12):
+        Number of blocks.
+    head_dim (`int`, *optional*, defaults to 64):
+        Width of one WKV head. `hidden_size` must be divisible by it.
+    num_heads (`int`, *optional*, defaults to 12):
+        Number of WKV heads; must equal `hidden_size // head_dim`.
+    decay_low_rank_dim (`int`, *optional*):
+        Rank of the decay (`w`) LoRA. When omitted, follows the RWKV-LM recipe:
+        `max(32, round(2.5 * sqrt(hidden_size) / 32) * 32)`.
+    a_low_rank_dim (`int`, *optional*):
+        Rank of the in-context-learning-rate (`a`) LoRA. Uses the same default
+        formula as `decay_low_rank_dim`.
+    v_low_rank_dim (`int`, *optional*):
+        Rank of the value-residual (`v`) LoRA. Defaults to
+        `max(32, round(1.7 * sqrt(hidden_size) / 32) * 32)`. Unused on layer 0,
+        which *produces* `v_first` instead of mixing towards it.
+    gate_low_rank_dim (`int`, *optional*):
+        Rank of the output-gate (`g`) LoRA. Defaults to
+        `max(32, round(5 * sqrt(hidden_size) / 32) * 32)`.
+    intermediate_size (`int`, *optional*):
+        Channel-mix inner width. Defaults to `4 * hidden_size`.
+    norm_eps (`float`, *optional*, defaults to 1e-05):
+        Epsilon of every LayerNorm/GroupNorm in the model.
+    norm_bias (`bool`, *optional*, defaults to `True`):
+        Whether the norms carry a bias.
+    tie_word_embeddings (`bool`, *optional*, defaults to `False`):
+        Whether to tie the input embedding and the LM head.
+    use_cache (`bool`, *optional*, defaults to `True`):
+        Whether to return the recurrent state.
+    wkv_state_dtype (`str`, *optional*, defaults to `"float32"`):
+        Precision the recurrent WKV state is carried and accumulated in,
+        independently of the activation dtype. The recurrence is unrolled over
+        the whole sequence, so a narrow state drifts; `"float32"` with fp16
+        activations is the combination the reference implementation uses.
+        `"float16"`/`"bfloat16"` trade that for a smaller state.
+    wkv_implementation (`str`, *optional*, defaults to `"eager"`):
+        Which WKV recurrence to use, by name, from
+        `models.rwkv7.modeling_rwkv7.RWKV7_WKV_FUNCTIONS`. `"eager"` is the
+        exact portable PyTorch path and preserves the reference token order for
+        both prefill and decode. `"chunked"` is an explicit faster
+        chunk-parallel rearrangement whose floating-point order differs. Register
+        an entry in that mapping to plug in a fused or varlen kernel without
+        forking the model.
+    bos_token_id (`int`, *optional*, defaults to 0):
+        Beginning-of-sequence id. The RWKV world tokenizer has no dedicated BOS
+        token and the reference implementation prepends nothing, so this exists to
+        satisfy `GenerationMixin` rather than to be emitted.
+    eos_token_id (`int`, *optional*, defaults to 0):
+        End-of-sequence id, id 0 in the RWKV world vocabulary.
+    pad_token_id (`int`, *optional*, defaults to 0):
+        Padding id, the same id 0. Set deliberately rather than left `None`:
+        `generate` needs one to pad a batch, and without it a batched call either
+        raised or fell back to the eos id with a warning on every step.
+
+    ```python
+    >>> from transformers import Rwkv7Config, Rwkv7Model
+
+    >>> configuration = Rwkv7Config()
+    >>> model = Rwkv7Model(configuration)
+    >>> configuration = model.config
+    ```"""
+
+    model_type = "rwkv7"
+    keys_to_ignore_at_inference = ["state"]
+
+    vocab_size: int = 65536
+    hidden_size: int = 768
+    num_hidden_layers: int = 12
+    head_dim: int = 64
+    num_heads: int = 12
+    decay_low_rank_dim: int | None = None
+    a_low_rank_dim: int | None = None
+    v_low_rank_dim: int | None = None
+    gate_low_rank_dim: int | None = None
+    # `None` rather than a number: a literal default is correct for the default
+    # `hidden_size` and silently wrong for every other one, so a config built as
+    # `Rwkv7Config(hidden_size=4096, num_heads=64)` would come back with a channel-mix
+    # four times narrower than the architecture it names. `__post_init__` resolves it,
+    # and the resolved value is written to `config.json` either way.
+    intermediate_size: int | None = None
+    norm_eps: float = 1e-5
+    norm_bias: bool = True
+    tie_word_embeddings: bool = False
+    use_cache: bool = True
+    wkv_state_dtype: str = "float32"
+    wkv_implementation: str = "eager"
+    bos_token_id: int | None = 0
+    eos_token_id: int | None = 0
+    pad_token_id: int | None = 0
+
+    def __post_init__(self, **kwargs):
+        rank_unit = self.hidden_size**0.5 / 32
+        if self.decay_low_rank_dim is None:
+            self.decay_low_rank_dim = max(32, round(2.5 * rank_unit) * 32)
+        if self.a_low_rank_dim is None:
+            self.a_low_rank_dim = max(32, round(2.5 * rank_unit) * 32)
+        if self.v_low_rank_dim is None:
+            self.v_low_rank_dim = max(32, round(1.7 * rank_unit) * 32)
+        if self.gate_low_rank_dim is None:
+            self.gate_low_rank_dim = max(32, round(5 * rank_unit) * 32)
+        if self.intermediate_size is None:
+            self.intermediate_size = 4 * self.hidden_size
+        if self.wkv_state_dtype not in ("float32", "float16", "bfloat16"):
+            raise ValueError(f"wkv_state_dtype must be float32/float16/bfloat16, got {self.wkv_state_dtype}")
+        if self.hidden_size % self.head_dim != 0:
+            raise ValueError(f"hidden_size {self.hidden_size} must be divisible by head_dim {self.head_dim}")
+        if self.num_heads != self.hidden_size // self.head_dim:
+            raise ValueError(
+                f"num_heads must be hidden_size // head_dim = {self.hidden_size // self.head_dim}, "
+                f"got {self.num_heads}"
+            )
+        super().__post_init__(**kwargs)
+
+
+__all__ = ["Rwkv7Config"]

generation_config.json

@@ -0,0 +1,6 @@
+{
+  "bos_token_id": 0,
+  "eos_token_id": 0,
+  "pad_token_id": 0,
+  "use_cache": true
+}

inference/generate.py

@@ -0,0 +1,239 @@
+from __future__ import annotations
+
+import argparse
+import importlib
+import os
+import re
+import sys
+from types import ModuleType
+from pathlib import Path
+from typing import Any
+
+import torch
+from transformers import StoppingCriteria, StoppingCriteriaList
+
+if __package__ in {None, ""}:
+    package_name = "_rwkv7_release_inference"
+    package = ModuleType(package_name)
+    package.__package__ = package_name
+    package.__path__ = [str(Path(__file__).resolve().parent)]
+    sys.modules[package_name] = package
+    load_model_and_tokenizer = importlib.import_module(
+        f"{package_name}.model_loader"
+    ).load_model_and_tokenizer
+else:
+    from .model_loader import load_model_and_tokenizer
+
+
+DTYPES = {
+    "bfloat16": torch.bfloat16,
+    "float16": torch.float16,
+    "float32": torch.float32,
+}
+STOP_TEXT = "\n\nUser:"
+THINK_RE = re.compile(r"\A<think>?\s*(.*?)\s*</think>?", re.DOTALL)
+
+
+class StopOnText(StoppingCriteria):
+    def __init__(self, tokenizer: Any, prompt_length: int, stop_text: str) -> None:
+        self.tokenizer = tokenizer
+        self.prompt_length = prompt_length
+        self.stop_text = stop_text
+        self.matched = False
+
+    def __call__(
+        self,
+        input_ids: torch.LongTensor,
+        scores: torch.FloatTensor,
+        **kwargs: Any,
+    ) -> bool:
+        del scores, kwargs
+        completion = self.tokenizer.decode(
+            input_ids[0, self.prompt_length :],
+            skip_special_tokens=False,
+        )
+        self.matched = self.stop_text in completion
+        return self.matched
+
+
+def _assistant_content(
+    completion: str, thinking: bool, *, close_incomplete: bool = False
+) -> str:
+    prefix = "<think" if thinking else "<think></think"
+    reply = prefix + completion
+    thinking_block = THINK_RE.match(reply)
+    if thinking:
+        if thinking_block is not None or not close_incomplete:
+            return reply.strip()
+        return f"{reply.rstrip()}\n</think>".strip()
+    return "" if thinking_block is None else reply[thinking_block.end() :].strip()
+
+
+def _prompt_ids(tokenizer: Any, messages: list[dict[str, str]], thinking: bool):
+    tokens = tokenizer.apply_chat_template(
+        messages,
+        tokenize=True,
+        add_generation_prompt=True,
+        thinking=thinking,
+        return_tensors="pt",
+    )
+    if hasattr(tokens, "input_ids"):
+        tokens = tokens.input_ids
+    elif isinstance(tokens, dict):
+        tokens = tokens["input_ids"]
+    if tokens.ndim == 1:
+        tokens = tokens.unsqueeze(0)
+    return tokens
+
+
+@torch.inference_mode()
+def generate_completion(
+    model: Any,
+    tokenizer: Any,
+    messages: list[dict[str, str]],
+    *,
+    device: str,
+    max_new_tokens: int,
+    temperature: float,
+    top_p: float,
+    thinking: bool,
+) -> str:
+    input_ids = _prompt_ids(tokenizer, messages, thinking).to(device)
+    prompt_length = input_ids.shape[1]
+    stop_on_text = StopOnText(tokenizer, prompt_length, STOP_TEXT)
+    generation: dict[str, Any] = {
+        "input_ids": input_ids,
+        "attention_mask": torch.ones_like(input_ids),
+        "max_new_tokens": max_new_tokens,
+        "do_sample": temperature > 0,
+        "eos_token_id": 0,
+        "pad_token_id": 0,
+        "stopping_criteria": StoppingCriteriaList([stop_on_text]),
+    }
+    if temperature > 0:
+        generation["temperature"] = temperature
+        generation["top_p"] = top_p
+    output = model.generate(**generation)
+    completion_ids = output[0, prompt_length:]
+    completion = tokenizer.decode(completion_ids, skip_special_tokens=True)
+    if STOP_TEXT in completion:
+        completion = completion.split(STOP_TEXT, 1)[0]
+    reached_token_limit = (
+        completion_ids.shape[0] >= max_new_tokens and not stop_on_text.matched
+    )
+    return _assistant_content(
+        completion,
+        thinking,
+        close_incomplete=reached_token_limit,
+    )
+
+
+def _interactive(
+    model: Any,
+    tokenizer: Any,
+    args: argparse.Namespace,
+) -> None:
+    messages: list[dict[str, str]] = []
+    print("RWKV-7 Goose — /clear resets the conversation, /exit quits.")
+    while True:
+        try:
+            prompt = input(">>> ")
+        except EOFError:
+            break
+        if prompt == "/exit":
+            break
+        if prompt == "/clear":
+            messages.clear()
+            continue
+        prompt = prompt.strip()
+        if not prompt:
+            continue
+        messages.append({"role": "user", "content": prompt})
+        completion = generate_completion(
+            model,
+            tokenizer,
+            messages,
+            device=args.device,
+            max_new_tokens=args.max_new_tokens,
+            temperature=args.temperature,
+            top_p=args.top_p,
+            thinking=args.thinking,
+        )
+        print(completion)
+        messages.append({"role": "assistant", "content": completion})
+
+
+def _file_prompts(
+    model: Any,
+    tokenizer: Any,
+    args: argparse.Namespace,
+) -> None:
+    text = Path(args.input_file).read_text(encoding="utf-8")
+    prompts = [prompt.strip() for prompt in re.split(r"\n\s*\n", text) if prompt.strip()]
+    if not prompts:
+        raise ValueError("input file contains no prompts")
+    for prompt in prompts:
+        completion = generate_completion(
+            model,
+            tokenizer,
+            [{"role": "user", "content": prompt}],
+            device=args.device,
+            max_new_tokens=args.max_new_tokens,
+            temperature=args.temperature,
+            top_p=args.top_p,
+            thinking=args.thinking,
+        )
+        print(f"Prompt: {prompt}")
+        print(f"Completion: {completion}")
+        print()
+
+
+def parse_args() -> argparse.Namespace:
+    parser = argparse.ArgumentParser(description="Generate with RWKV-7 Goose")
+    parser.add_argument("--model", required=True, help="Hub repo ID or local model directory")
+    mode = parser.add_mutually_exclusive_group(required=True)
+    mode.add_argument("--interactive", action="store_true")
+    mode.add_argument("--input-file")
+    parser.add_argument("--device", default="cuda")
+    parser.add_argument("--dtype", choices=("auto", *DTYPES), default="auto")
+    parser.add_argument("--state-dtype", choices=DTYPES, default="float32")
+    parser.add_argument("--backend", choices=("auto", "torch", "tilelang"), default="auto")
+    parser.add_argument("--max-new-tokens", type=int, default=300)
+    parser.add_argument("--temperature", type=float, default=1.0)
+    parser.add_argument("--top-p", type=float, default=0.5)
+    parser.add_argument("--seed", type=int, default=33377335)
+    parser.add_argument("--thinking", action="store_true")
+    return parser.parse_args()
+
+
+def main() -> None:
+    args = parse_args()
+    if int(os.getenv("WORLD_SIZE", "1")) != 1:
+        raise RuntimeError("the bundled runtime supports one process and one GPU")
+    if int(os.getenv("RANK", "0")) != 0 or int(os.getenv("LOCAL_RANK", "0")) != 0:
+        raise RuntimeError("RANK and LOCAL_RANK must be zero")
+    if args.max_new_tokens <= 0:
+        raise ValueError("max-new-tokens must be positive")
+    if args.temperature < 0:
+        raise ValueError("temperature must be non-negative")
+    if not 0 < args.top_p <= 1:
+        raise ValueError("top-p must be in (0, 1]")
+    torch.manual_seed(args.seed)
+    model, tokenizer = load_model_and_tokenizer(
+        args.model,
+        device=args.device,
+        dtype=None if args.dtype == "auto" else DTYPES[args.dtype],
+        backend=args.backend,
+        state_dtype=args.state_dtype,
+    )
+    model.set_kernel_backend(args.backend)
+    if args.backend == "tilelang":
+        model.prepare_inference_weights()
+    if args.interactive:
+        _interactive(model, tokenizer, args)
+    else:
+        _file_prompts(model, tokenizer, args)
+
+
+if __name__ == "__main__":
+    main()

inference/model_loader.py

@@ -0,0 +1,178 @@
+from __future__ import annotations
+
+import json
+from pathlib import Path
+
+import torch
+from huggingface_hub import snapshot_download
+from safetensors.torch import load_file
+from transformers import AutoTokenizer, PreTrainedConfig
+
+from .runtime import RWKV7Config, RWKV7ForCausalLM
+
+
+SQUEEZE_MARKERS = (
+    ".x_",
+    ".k_",
+    "att.r",
+    "att.w",
+    "att.v0",
+    "att.v1",
+    "att.v2",
+    "att.a",
+    "att.g",
+)
+
+
+def _resolve_model(model: str) -> tuple[Path, bool]:
+    local = Path(model).expanduser()
+    if local.is_dir():
+        return local.resolve(), False
+    return Path(snapshot_download(model)), True
+
+
+def _optimized_config(native: dict, backend: str, state_dtype: str) -> RWKV7Config:
+    return RWKV7Config(
+        vocab_size=native["vocab_size"],
+        hidden_size=native["hidden_size"],
+        num_hidden_layers=native["num_hidden_layers"],
+        head_size=native["head_dim"],
+        intermediate_size=native["intermediate_size"],
+        decay_lora_rank=native["decay_low_rank_dim"],
+        a_lora_rank=native["a_low_rank_dim"],
+        gate_lora_rank=native["gate_low_rank_dim"],
+        value_lora_rank=native["v_low_rank_dim"],
+        layer_norm_epsilon=native.get("norm_eps", 1e-5),
+        use_cache=native.get("use_cache", True),
+        kernel_backend=backend,
+        recurrent_state_dtype=state_dtype,
+        tie_word_embeddings=native.get("tie_word_embeddings", False),
+        bos_token_id=native.get("bos_token_id"),
+        eos_token_id=native.get("eos_token_id", 0),
+        pad_token_id=native.get("pad_token_id", 0),
+    )
+
+
+def _native_key_to_optimized(key: str) -> str:
+    return key.removeprefix("rwkv7.")
+
+
+def _native_tensor_to_optimized(key: str, tensor: torch.Tensor) -> torch.Tensor:
+    return tensor.squeeze() if any(marker in key for marker in SQUEEZE_MARKERS) else tensor
+
+
+def _checked_weight(path: Path, *, hub_snapshot: bool) -> Path:
+    if path.is_symlink():
+        if not hub_snapshot:
+            raise RuntimeError(f"local safetensor must not be a symlink: {path.name}")
+        path = path.resolve(strict=True)
+    if not path.is_file():
+        raise RuntimeError(f"safetensor must be a regular file: {path.name}")
+    return path
+
+
+def _weight_plan(
+    model_dir: Path, *, hub_snapshot: bool
+) -> list[tuple[Path, set[str] | None]]:
+    present = sorted(model_dir.glob("model*.safetensors"))
+    if not present:
+        raise FileNotFoundError(f"no safetensors found in {model_dir}")
+    checked = {
+        path.name: _checked_weight(path, hub_snapshot=hub_snapshot) for path in present
+    }
+    index_path = model_dir / "model.safetensors.index.json"
+    if not index_path.is_file():
+        if [path.name for path in present] != ["model.safetensors"]:
+            raise RuntimeError("multiple safetensors require model.safetensors.index.json")
+        return [(checked["model.safetensors"], None)]
+
+    index = json.loads(index_path.read_text(encoding="utf-8"))
+    weight_map = index.get("weight_map")
+    if not isinstance(weight_map, dict) or not weight_map:
+        raise RuntimeError("model.safetensors.index.json has no weight_map")
+    keys_by_file: dict[str, set[str]] = {}
+    for key, filename in weight_map.items():
+        if not isinstance(key, str) or not isinstance(filename, str):
+            raise RuntimeError("invalid weight_map entry")
+        keys_by_file.setdefault(filename, set()).add(key)
+    if set(keys_by_file) != {path.name for path in present}:
+        raise RuntimeError("indexed and present safetensors files differ")
+    return [(checked[filename], keys_by_file[filename]) for filename in sorted(keys_by_file)]
+
+
+def _load_native_model(
+    model_dir: Path,
+    native_config: dict,
+    backend: str,
+    state_dtype: str,
+    *,
+    hub_snapshot: bool,
+) -> RWKV7ForCausalLM:
+    config = _optimized_config(native_config, backend, state_dtype)
+    with torch.device("meta"):
+        model = RWKV7ForCausalLM(config)
+    expected = set(model.state_dict())
+    seen: set[str] = set()
+    for weight_file, indexed_keys in _weight_plan(
+        model_dir, hub_snapshot=hub_snapshot
+    ):
+        native_shard = load_file(weight_file, device="cpu")
+        if indexed_keys is not None and set(native_shard) != indexed_keys:
+            raise RuntimeError(f"tensor keys in {weight_file.name} do not match the index")
+        shard = {}
+        for native_key, tensor in native_shard.items():
+            optimized_key = _native_key_to_optimized(native_key)
+            if optimized_key in seen or optimized_key in shard:
+                raise RuntimeError(f"duplicate optimized tensor key: {optimized_key}")
+            shard[optimized_key] = _native_tensor_to_optimized(optimized_key, tensor)
+        unexpected = sorted(set(shard) - expected)
+        if unexpected:
+            raise RuntimeError(f"native checkpoint has unexpected keys: {unexpected}")
+        model.load_state_dict(shard, strict=False, assign=True)
+        seen.update(shard)
+    missing = sorted(expected - seen)
+    if missing:
+        raise RuntimeError(f"native checkpoint is missing optimized runtime keys: {missing}")
+    return model
+
+
+def _load_tokenizer(model_dir: Path):
+    return AutoTokenizer.from_pretrained(
+        model_dir,
+        config=PreTrainedConfig(),
+        local_files_only=True,
+    )
+
+
+def load_model_and_tokenizer(
+    model: str,
+    *,
+    device: str,
+    dtype: torch.dtype | None,
+    backend: str,
+    state_dtype: str,
+):
+    model_dir, hub_snapshot = _resolve_model(model)
+    native_config = json.loads((model_dir / "config.json").read_text(encoding="utf-8"))
+    architectures = set(native_config.get("architectures", []))
+    if architectures != {"Rwkv7ForCausalLM"}:
+        raise ValueError(f"unsupported RWKV-7 architecture: {sorted(architectures)}")
+    loaded = _load_native_model(
+        model_dir,
+        native_config,
+        backend,
+        state_dtype,
+        hub_snapshot=hub_snapshot,
+    )
+    if dtype is None:
+        dtype_name = str(native_config.get("dtype", "bfloat16")).removeprefix("torch.")
+        try:
+            dtype = {
+                "bfloat16": torch.bfloat16,
+                "float16": torch.float16,
+                "float32": torch.float32,
+            }[dtype_name]
+        except KeyError as error:
+            raise ValueError(f"unsupported model dtype: {dtype_name}") from error
+    loaded = loaded.to(device=device, dtype=dtype).eval()
+    return loaded, _load_tokenizer(model_dir)

inference/requirements.txt

@@ -0,0 +1,5 @@
+transformers>=5.15,<6
+huggingface-hub>=0.34
+safetensors>=0.5
+jinja2>=3.1,<4
+tilelang==0.1.12

model.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:84ccbb857c84e00cefc48b233937ada79c411e491df25fb21aed23237f39a14f
+size 3055418240

modeling_rwkv7.py

@@ -0,0 +1,1186 @@
+# Copyright 2026 The RWKV team and The HuggingFace Inc. team. All rights reserved.
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+#     http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+"""PyTorch RWKV-7 ("Goose") model."""
+
+import math
+from dataclasses import dataclass
+
+import torch
+from torch import nn
+
+from transformers import initialization as init
+from transformers.cache_utils import Cache, LinearAttentionLayer
+from transformers.generation import GenerationMixin
+from transformers.modeling_layers import GradientCheckpointingLayer
+from transformers.modeling_utils import PreTrainedModel
+from transformers.utils import ModelOutput, auto_docstring, can_return_tuple, logging
+from .configuration_rwkv7 import Rwkv7Config
+
+
+logger = logging.get_logger(__name__)
+
+
+def _rwkv7_lora_orthogonal_(tensor: torch.Tensor, scale: float = 0.1) -> None:
+    """RWKV-LM compensates the gain when a low-rank output factor is tall."""
+    rows, columns = tensor.shape
+    gain = math.sqrt(rows / columns) if rows > columns else 1.0
+    init.orthogonal_(tensor, gain=gain * scale)
+
+
+class Rwkv7TokenShift(nn.Module):
+    """`prev_token(x)`: the previous token's hidden state, zero at sequence start.
+
+    Carries one vector per layer per stream across forward calls, which is what
+    makes incremental decoding exact rather than approximate.
+    """
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        shift_state: torch.Tensor | None,
+        cu_seq_lens: torch.Tensor | None = None,
+        keep: torch.Tensor | None = None,
+    ):
+        # x: [batch, seq_len, hidden]; shift_state: [batch, hidden] or None
+        if shift_state is None:
+            prev = torch.zeros_like(x[:, :1])
+        else:
+            prev = shift_state[:, None]
+        shifted = torch.cat([prev, x[:, :-1]], dim=1)
+        if cu_seq_lens is not None:
+            # In a packed row the token before a segment's first one belongs to the
+            # PREVIOUS sequence. Resetting the recurrent state per segment does not
+            # cover this (the shift reaches back through it), so the first token of
+            # each segment gets the zero shift a sequence start is supposed to see.
+            positions = torch.arange(x.shape[1], device=x.device)
+            starts = (positions[:, None] == cu_seq_lens[None, :-1]).any(dim=1)
+            shifted = torch.where(starts[None, :, None], torch.zeros_like(shifted), shifted)
+        if keep is None:
+            return shifted, x[:, -1]
+        # The state handed back must be the last REAL token, not the last position.
+        # They coincide under left padding, which is why this went unnoticed; under
+        # right padding `x[:, -1]` is a blanked pad, so a continuation resumes from
+        # zero instead of from where the sequence actually got to.
+        # Integer arithmetic, not the activation dtype. `keep` arrives cast to the
+        # model's dtype, and bf16 carries eight mantissa bits: past position 256 the
+        # products are no longer distinct, `argmax` returns the first of a tie, and
+        # the state comes back from a token several places short of the last real one
+        # -- 1022 for 1023 at length 1024, 4088 for 4095 at 4096. An all-ones mask is
+        # enough to trigger it, which is what `generate` sends, so this was not
+        # confined to padded batches.
+        positions = torch.arange(x.shape[1], device=x.device)
+        last_real = ((keep.squeeze(-1) > 0) * positions).argmax(dim=-1)
+        return shifted, x[torch.arange(x.shape[0], device=x.device), last_real]
+
+
+# e^-0.5. The decay LoRA emits w_log = -INV_SQRT_E * sigmoid(...), so w_log lies in
+# (-e^-0.5, 0) and the per-step decay exp(w_log) lies in (exp(-e^-0.5), 1), i.e.
+# (0.5452, 1). See RWKV-7 reference. Note the floor is exp(-e^-0.5), not e^-0.5: this
+# comment said the latter, and `rwkv7_chunked` built its chunk_size bound on it.
+_INV_SQRT_E = 0.6065306597126334
+
+
+def _rwkv7_matmul(x: torch.Tensor, weight: torch.Tensor, tokenwise: bool) -> torch.Tensor:
+    """Preserve decode GEMM shapes across a sequence when exact execution is requested."""
+    if not tokenwise or x.ndim < 3 or x.shape[1] == 1:
+        return x @ weight
+    return torch.stack([x[:, token_index] @ weight for token_index in range(x.shape[1])], dim=1)
+
+
+def _rwkv7_linear(module: nn.Linear, x: torch.Tensor, tokenwise: bool) -> torch.Tensor:
+    if not tokenwise or x.ndim < 3 or x.shape[1] == 1:
+        return module(x)
+    return torch.stack([module(x[:, token_index]) for token_index in range(x.shape[1])], dim=1)
+
+
+def rwkv7_recurrent(
+    r: torch.Tensor,
+    w_log: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    kk: torch.Tensor,
+    a: torch.Tensor,
+    state: torch.Tensor,
+    compute_dtype: torch.dtype = torch.float32,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Reference RWKV-7 WKV recurrence (generalised delta rule).
+
+    All inputs are `[batch, seq_len, num_heads, head_dim]`; `state` is
+    `[batch, num_heads, head_dim, head_dim]` and is carried across calls.
+
+    Per token, with `S` the per-head state::
+
+        decay = exp(w_log)
+        vk    = v @ k
+        ab    = (-kk) @ (kk * a)
+        S     = S * decay + S @ ab + vk
+        out   = S @ r
+
+    The state axes are (value, key), as in the reference implementation. Products
+    made from activations keep the activation dtype; the carried-state arithmetic
+    uses `compute_dtype`, and a narrower configured state is stored at every token
+    boundary before the output projection. Those boundaries are observable in a
+    recurrent BF16 model, so an algebraically transposed FP32 implementation is not
+    equivalent.
+    """
+    batch, seq_len, num_heads, head_dim = r.shape
+    dtype = r.dtype
+    state_dtype = state.dtype
+    state = state.to(compute_dtype)
+    out = torch.empty(batch, seq_len, num_heads, head_dim, device=r.device, dtype=compute_dtype)
+
+    for t in range(seq_len):
+        decay = torch.exp(w_log[:, t])[:, :, None, :]  # [B, H, 1, K]
+        kk_t = kk[:, t]
+        value_key = v[:, t][..., None] @ k[:, t][:, :, None, :]
+        anti_value_key = (-kk_t)[..., None] @ (kk_t * a[:, t])[:, :, None, :]
+        matrix = state.to(compute_dtype)
+        state = matrix * decay.to(compute_dtype)
+        state = state + matrix @ anti_value_key.to(compute_dtype)
+        state = state + value_key.to(compute_dtype)
+        if state_dtype != compute_dtype:
+            state = state.to(state_dtype)
+        out[:, t] = (state.to(dtype) @ r[:, t][..., None]).squeeze(-1).to(compute_dtype)
+
+    return out.to(dtype), state
+
+
+def rwkv7_chunked(
+    r: torch.Tensor,
+    w_log: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    kk: torch.Tensor,
+    a: torch.Tensor,
+    state: torch.Tensor,
+    chunk_size: int = 64,
+    compute_dtype: torch.dtype = torch.float32,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Chunk-parallel form of [`rwkv7_recurrent`], for sequences.
+
+    The step is `S_t = A_t S_{t-1} + k_t v_t^T` with `A_t = diag(w_t) - b_t kk_t^T`
+    and `b_t = kk_t * a_t`, i.e. diagonal-plus-rank-one. Substituting
+    `S_t = diag(c_t) P_t` with the running decay `c_t = prod_{s<=t} w_s` removes the
+    diagonal part and leaves a plain delta rule::
+
+        P_t = (I - b~_t q~_t^T) P_{t-1} + k~_t v_t^T
+        b~ = b / c,  k~ = k / c,  q~ = kk * c_{t-1},  r~ = r * c
+
+    Writing `u_t = q~_t^T P_{t-1}` turns a whole chunk into one unit-lower-triangular
+    system, so the chunk is a handful of matmuls instead of `chunk_size` sequential
+    steps; only the chunk-to-chunk carry stays serial::
+
+        (I + tril(Q~ B~^T, -1)) U = Q~ P_0 + tril(Q~ K~^T, -1) V
+        O   = R~ P_0 + tril(R~ K~^T, 0) V - tril(R~ B~^T, 0) U
+        P_C = P_0 + K~^T V - B~^T U
+
+    `chunk_size` is bounded by that division by `c`, but by OVERFLOW rather than by
+    precision, and the difference is worth a factor of four. The per-step decay is at
+    least `exp(-e^-0.5)` = 0.5452, so at the worst case, every channel pinned at that
+    floor, `1/c` grows like `e^(e^-0.5 * chunk_size)`. That reaches 7.2e16 at 64 and
+    fp32 tops out near 3.4e38, so there are twenty-odd decades of headroom left at the
+    default, and the ceiling is `ln(finfo.max) / e^-0.5` = 146 in fp32.
+
+    That derivation is worth stating carefully because the first version of it was
+    wrong in a way the numbers hid. It said the decay floor was `e^-0.5` rather than
+    `exp(-e^-0.5)`, hence growth like `e^(0.5 * chunk_size)`, which puts the ceiling at
+    177, and measured at the decay floor this function returns all-NaN from 147 up. The
+    quoted 7.2e16 came from the correct law all along (`e^(0.5*64)` is 7.9e13), so the
+    arithmetic had been done right and written up wrong, and following the prose rather
+    than the number led into the overflow band. `chunk_size` is checked below rather
+    than left to that reasoning.
+
+    Precision does not degrade with it, because the substitution is a similarity
+    transform: whatever `1/c` inflates, `c` deflates again on the way out, and a
+    common scale factor does not move a floating-point relative error. Measured at
+    the decay floor, `T=256`, against the sequential form: chunk 16 gives 2.0e-07 and
+    chunk 64 gives 2.7e-07, both fp32 noise, while the recurrence runs 3.9x faster
+    (3.771 ms -> 0.913 ms). 16 was the conservative reading of the same bound.
+    """
+    batch, seq_len, num_heads, head_dim = r.shape
+    dtype = r.dtype
+    r, w_log, k, v, kk, a = (t.to(compute_dtype) for t in (r, w_log, k, v, kk, a))
+    state_dtype = state.dtype
+    # The chunk derivation below uses the transposed (key, value) form internally.
+    state = state.transpose(-2, -1).to(compute_dtype)
+    # Everything that does not touch the carried state is computed for ALL chunks at
+    # once. Only the state recurrence is serial, and it was pulling the rest of the
+    # arithmetic into the Python loop with it: at T=256 and chunk 64 that was four
+    # iterations of ten-odd launches each, plus three constant matrices rebuilt every
+    # iteration. The loop below now does the state-dependent terms and nothing else.
+    # Never pad past the sequence: a 16-token prefill grouped into 64-wide chunks does
+    # four times the arithmetic for the same answer, which cost 1x16 nine points before
+    # this line existed. The old loop avoided it by shortening its last chunk.
+    chunk_size = min(chunk_size, seq_len)
+    # Refuse rather than return NaN. `1/c` reaches `e^(e^-0.5 * chunk_size)` at the
+    # decay floor, so the widest chunk this dtype can carry is derived from the dtype
+    # rather than written down: 146 in fp32, 12 in fp16. A caller who raises
+    # `chunk_size` for a longer prefill gets an error naming the limit instead of
+    # all-NaN logits several layers later, which is what the previous version did.
+    widest = int(math.log(torch.finfo(compute_dtype).max) / _INV_SQRT_E)
+    if chunk_size > widest:
+        raise ValueError(
+            f"chunk_size={chunk_size} overflows {compute_dtype}: the running decay's "
+            f"reciprocal grows like e^(e^-0.5 * chunk_size), so the widest chunk that "
+            f"stays finite is {widest}. Lower chunk_size, or raise compute_dtype."
+        )
+    chunks = (seq_len + chunk_size - 1) // chunk_size
+    padded = chunks * chunk_size
+    if padded != seq_len:
+        pad = (0, 0, 0, 0, 0, padded - seq_len)
+        r, k, v, kk, a = (torch.nn.functional.pad(t, pad) for t in (r, k, v, kk, a))
+        # A pad step must be the identity for the recurrence: decay 1, nothing added.
+        w_log = torch.nn.functional.pad(w_log, pad)
+    grouped = lambda t: t.reshape(batch, chunks, chunk_size, num_heads, head_dim)  # noqa: E731
+    rg, kg, vg, kkg, ag, wg = (grouped(t) for t in (r, k, v, kk, a, w_log))
+
+    w_c = torch.exp(wg)
+    c = torch.cumprod(w_c, dim=2)
+    c_prev = c / w_c
+    bg = kkg * ag
+    b_t, k_t = bg / c, kg / c
+    q_t, r_t = kkg * c_prev, rg * c
+
+    span = chunk_size
+    tri = torch.ones(span, span, device=r.device, dtype=r.dtype).tril(-1)
+    causal = torch.ones(span, span, device=r.device, dtype=r.dtype).tril(0)
+    eye = torch.eye(span, device=r.device, dtype=r.dtype)
+
+    # [batch, chunks, heads, span, span] -- one launch each instead of one per chunk.
+    qb = torch.einsum("bcthn,bcshn->bchts", q_t, b_t) * tri
+    qk = torch.einsum("bcthn,bcshn->bchts", q_t, k_t) * tri
+    rk = torch.einsum("bcthn,bcshn->bchts", r_t, k_t) * causal
+    rb = torch.einsum("bcthn,bcshn->bchts", r_t, b_t) * causal
+    lhs = eye + qb
+    qkv = torch.einsum("bchts,bcshv->bchtv", qk, vg)
+    rkv = torch.einsum("bchts,bcshv->bchtv", rk, vg)
+    c_last = c[:, :, -1].unsqueeze(-1)
+
+    # `k~^T v` carries no state either, so it joins the batched half above; only the
+    # two products that read the carried state and the solve that depends on them are
+    # genuinely serial.
+    kv = torch.einsum("bcthn,bcthv->bchnv", k_t, vg)
+    # Both state products read the SAME state, so they are one matmul over a stacked
+    # token axis rather than two -- the loop is short and launch-bound, and this is one
+    # launch per chunk instead of two.
+    qr = torch.cat([q_t, r_t], dim=2)
+
+    outputs = []
+    for i in range(chunks):
+        qr_s = torch.einsum("bthn,bhnv->bhtv", qr[:, i], state)
+        rhs = qr_s[:, :, :chunk_size] + qkv[:, i]
+        u = torch.linalg.solve_triangular(lhs[:, i], rhs, upper=False, unitriangular=True)
+        out_c = qr_s[:, :, chunk_size:] + rkv[:, i] - torch.einsum("bhts,bhsv->bhtv", rb[:, i], u)
+        outputs.append(out_c.permute(0, 2, 1, 3))
+        state = c_last[:, i] * (state + kv[:, i] - torch.einsum("bthn,bhtv->bhnv", b_t[:, i], u))
+
+    state = state.transpose(-2, -1).contiguous().to(state_dtype)
+    return torch.cat(outputs, dim=1)[:, :seq_len].to(dtype), state
+
+
+def rwkv7_eager(
+    r: torch.Tensor,
+    w_log: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    kk: torch.Tensor,
+    a: torch.Tensor,
+    state: torch.Tensor,
+    cu_seq_lens: torch.Tensor | None = None,
+    **kwargs,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Portable exact WKV, evaluated in reference token order.
+
+    With `cu_seq_lens` the row is a *packed* batch: several independent sequences
+    laid end to end, the layout a varlen kernel consumes. Each one has to start
+    from a fresh state, so they are run in turn and their outputs concatenated;
+    the state returned is the last segment's, which is the one a continuation
+    would resume from. This is the reference behaviour, not the fast path: a
+    fused varlen kernel does the same segments in one launch, which is exactly
+    the kind of implementation [`RWKV7_WKV_FUNCTIONS`] exists to let you drop in.
+
+    This default deliberately stays sequential for a multi-token prefill. The
+    chunk-parallel rearrangement is much faster, but changes floating-point operation
+    order and is therefore available only through the explicit `"chunked"` registry
+    entry. Automatic execution must remain the reference users can compare it to.
+    """
+    if cu_seq_lens is not None:
+        bounds = cu_seq_lens.tolist()
+        outputs = []
+        for start, stop in zip(bounds[:-1], bounds[1:]):
+            if stop <= start:
+                continue
+            segment_state = torch.zeros_like(state)
+            out, state = rwkv7_eager(
+                r[:, start:stop],
+                w_log[:, start:stop],
+                k[:, start:stop],
+                v[:, start:stop],
+                kk[:, start:stop],
+                a[:, start:stop],
+                segment_state,
+            )
+            outputs.append(out)
+        return torch.cat(outputs, dim=1), state
+
+    return rwkv7_recurrent(r, w_log, k, v, kk, a, state)
+
+
+def rwkv7_parallel(
+    r: torch.Tensor,
+    w_log: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    kk: torch.Tensor,
+    a: torch.Tensor,
+    state: torch.Tensor,
+    cu_seq_lens: torch.Tensor | None = None,
+    **kwargs,
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Explicit chunk-parallel WKV with the same packed-sequence contract as eager."""
+    if cu_seq_lens is not None:
+        bounds = cu_seq_lens.tolist()
+        outputs = []
+        for start, stop in zip(bounds[:-1], bounds[1:]):
+            if stop <= start:
+                continue
+            out, state = rwkv7_chunked(
+                r[:, start:stop],
+                w_log[:, start:stop],
+                k[:, start:stop],
+                v[:, start:stop],
+                kk[:, start:stop],
+                a[:, start:stop],
+                torch.zeros_like(state),
+            )
+            outputs.append(out)
+        return torch.cat(outputs, dim=1), state
+    return rwkv7_chunked(r, w_log, k, v, kk, a, state)
+
+
+RWKV7_WKV_FUNCTIONS = {"eager": rwkv7_eager, "chunked": rwkv7_parallel}
+
+
+class Rwkv7Attention(nn.Module):
+    """RWKV-7 time-mixing block (the recurrent replacement for self-attention)."""
+
+    def __init__(self, config: Rwkv7Config, layer_id: int):
+        super().__init__()
+        self.config = config
+        self.layer_id = layer_id
+        C = config.hidden_size
+        self.hidden_size = C
+        self.num_heads = config.num_heads
+        self.head_dim = config.head_dim
+        self.state_dtype = getattr(torch, config.wkv_state_dtype)
+
+        self.time_shift = Rwkv7TokenShift()
+
+        # per-channel token-shift mixes (kept at Bo's (1, 1, C) shape)
+        self.x_r = nn.Parameter(torch.zeros(1, 1, C))
+        self.x_w = nn.Parameter(torch.zeros(1, 1, C))
+        self.x_k = nn.Parameter(torch.zeros(1, 1, C))
+        self.x_v = nn.Parameter(torch.zeros(1, 1, C))
+        self.x_a = nn.Parameter(torch.zeros(1, 1, C))
+        self.x_g = nn.Parameter(torch.zeros(1, 1, C))
+
+        # LoRA factors as raw tensors, exactly as the reference stores them
+        self.w1 = nn.Parameter(torch.zeros(C, config.decay_low_rank_dim))
+        self.w2 = nn.Parameter(torch.zeros(config.decay_low_rank_dim, C))
+        self.w0 = nn.Parameter(torch.zeros(1, 1, C))
+        self.a1 = nn.Parameter(torch.zeros(C, config.a_low_rank_dim))
+        self.a2 = nn.Parameter(torch.zeros(config.a_low_rank_dim, C))
+        self.a0 = nn.Parameter(torch.zeros(1, 1, C))
+        self.g1 = nn.Parameter(torch.zeros(C, config.gate_low_rank_dim))
+        self.g2 = nn.Parameter(torch.zeros(config.gate_low_rank_dim, C))
+        # The value-residual LoRA exists on every layer in a reference checkpoint,
+        # but layer 0 PRODUCES v_first instead of mixing towards it, so its copy is
+        # never read. It is registered anyway so that loading is lossless.
+        self.v1 = nn.Parameter(torch.zeros(C, config.v_low_rank_dim))
+        self.v2 = nn.Parameter(torch.zeros(config.v_low_rank_dim, C))
+        self.v0 = nn.Parameter(torch.zeros(1, 1, C))
+
+        self.k_k = nn.Parameter(torch.zeros(1, 1, C))
+        self.k_a = nn.Parameter(torch.zeros(1, 1, C))
+        self.r_k = nn.Parameter(torch.zeros(config.num_heads, config.head_dim))
+
+        self.receptance = nn.Linear(C, C, bias=False)
+        self.key = nn.Linear(C, C, bias=False)
+        self.value = nn.Linear(C, C, bias=False)
+        self.output = nn.Linear(C, C, bias=False)
+        # GroupNorm over heads, matching the reference's per-head normalisation.
+        #
+        # The reference hardcodes `eps=64e-5`, which is `head_dim * 1e-5` at its
+        # head_dim of 64 -- NOT `num_heads * 1e-5`. The two coincide only when a
+        # model happens to have as many heads as channels per head, which for
+        # head_dim 64 means hidden_size 4096 exactly. Scaling by `num_heads` is
+        # therefore right on the 7.2B and wrong everywhere else: 2x low on the 1.5B,
+        # 5.33x low on the 0.1B. It is written against `head_dim` because that is
+        # the axis GroupNorm actually reduces over, so the constant tracks the
+        # reference at any width.
+        self.ln_x = nn.GroupNorm(config.num_heads, C, eps=config.norm_eps * config.head_dim)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        v_first: torch.Tensor | None,
+        shift_state: torch.Tensor | None,
+        wkv_state: torch.Tensor | None,
+        keep: torch.Tensor | None = None,
+        cu_seq_lens: torch.Tensor | None = None,
+    ):
+        batch, seq_len, C = hidden_states.shape
+        H, N = self.num_heads, self.head_dim
+        tokenwise = self.config.wkv_implementation != "chunked"
+
+        shifted, new_shift_state = self.time_shift(hidden_states, shift_state, cu_seq_lens, keep)
+        delta = shifted - hidden_states
+        xr = hidden_states + self.x_r * delta
+        xw = hidden_states + self.x_w * delta
+        xk = hidden_states + self.x_k * delta
+        xv = hidden_states + self.x_v * delta
+        xa = hidden_states + self.x_a * delta
+        xg = hidden_states + self.x_g * delta
+
+        r = _rwkv7_linear(self.receptance, xr, tokenwise)
+        k = _rwkv7_linear(self.key, xk, tokenwise)
+        v = _rwkv7_linear(self.value, xv, tokenwise)
+
+        w_log, a, g, v_gate = self.lora_gates(xw, xa, xg, None if self.layer_id == 0 else xv, tokenwise)
+
+        if keep is not None:
+            # A padding position has to leave the recurrent state exactly as it
+            # found it, which takes three things and not one.
+            #
+            # The decay is held at w = exp(0) = 1, so the transition is the identity.
+            #
+            # `k` and `v` are zeroed EXPLICITLY rather than relying on the blanked
+            # hidden state to make them zero: the projections are bias-free, so that
+            # only holds for a pad with nothing before it.
+            # A pad that FOLLOWS a real token still receives that token's hidden
+            # state through the token shift, so `delta = shifted - 0` is non-zero,
+            # and so are `k` and `v`. The update term `k v^T` then entered the state
+            # on every right-padded batch and on any left-padded batch continued
+            # from a carried shift state.
+            w_log = w_log * keep
+            k = k * keep
+            v = v * keep
+
+        if self.layer_id == 0:
+            v_first = v
+        else:
+            v = v + (v_first - v) * v_gate
+
+        kk = k * self.k_k
+        kk = torch.nn.functional.normalize(kk.view(batch, seq_len, H, N), dim=-1, p=2.0).view(batch, seq_len, C)
+        if keep is not None:
+            # A blanked padding position makes `k` exactly zero, so this normalises a
+            # zero vector. `F.normalize` divides by `max(norm, 1e-12)`, and 1e-12 is
+            # below the smallest fp16 subnormal -- so in fp16 the divisor really is
+            # zero and every padded row comes out NaN. (In fp32 it is representable
+            # and the whole thing looks fine, which is why a test on a small fp32
+            # model missed it entirely.) `where` is used rather than a multiply
+            # because NaN * 0 is still NaN.
+            kk = torch.where(keep.bool().expand_as(kk), kk, torch.zeros_like(kk))
+        k = k * (1 + (a - 1) * self.k_a)
+
+        if wkv_state is None:
+            wkv_state = torch.zeros(batch, H, N, N, device=r.device, dtype=self.state_dtype)
+
+        def _heads(t):
+            return t.view(batch, seq_len, H, N)
+
+        # Named rather than indexed: an unregistered key would otherwise surface as a
+        # bare `KeyError: 'chunked'` from inside the forward, several frames from the
+        # config field that caused it, and a caller looping over shapes would record it
+        # as "this shape did not run" instead of "this model was never built". The
+        # registry is open by design, so this cannot be validated in the config.
+        try:
+            wkv = RWKV7_WKV_FUNCTIONS[self.config.wkv_implementation]
+        except KeyError:
+            raise ValueError(
+                f"wkv_implementation={self.config.wkv_implementation!r} is not registered. "
+                f"Known: {sorted(RWKV7_WKV_FUNCTIONS)}. Register your own with "
+                "`RWKV7_WKV_FUNCTIONS['name'] = fn` before building the model."
+            ) from None
+        y, wkv_state = wkv(
+            _heads(r),
+            _heads(w_log),
+            _heads(k),
+            _heads(v),
+            _heads(kk),
+            _heads(a),
+            wkv_state,
+            cu_seq_lens=cu_seq_lens,
+        )
+
+        # `reshape`, not `view`: what comes back from the WKV is whatever layout that
+        # implementation produced, and a registered kernel -- or inductor, which is
+        # free to pick its own -- can hand back a `[batch, seq_len, heads, head_dim]`
+        # that is strided as if it were `[batch, heads, seq_len, head_dim]`. `view`
+        # then raises, and it raises only when batch and seq_len are BOTH greater
+        # than one, which is why a compiled 16x16 forward failed while 1xT and Bx1
+        # both passed.
+        y = self.ln_x(y.reshape(batch * seq_len, C)).view(batch, seq_len, C)
+        # r·k·r_k summed per head, broadcast back over the head's value channels
+        bonus = ((_heads(r) * _heads(k) * self.r_k).sum(dim=-1, keepdim=True) * _heads(v)).reshape(batch, seq_len, C)
+        y = _rwkv7_linear(self.output, (y + bonus) * g, tokenwise)
+        return y, v_first, new_shift_state, wkv_state
+
+    def lora_gates(self, xw, xa, xg, xv, tokenwise):
+        """The four low-rank chains, as one unit.
+
+        Each is a rank-r down projection, an activation, and an up projection, and
+        they differ only in where the activation sits: `w` has `tanh` between the
+        two, `g` has `sigmoid` between them and nothing after, `a` and `v` have
+        nothing between and `sigmoid` after. Together they are a dozen tiny
+        matrix-vector products per layer whose weights are a rounding error of the
+        model's bytes -- so at batch 1 they cost latency, not bandwidth, which is
+        the shape a fused kernel improves and a portable implementation cannot.
+
+        Kept as one method for exactly that reason: it is the unit worth replacing,
+        and replacing it should not need a fork. `xv` is None on layer 0, which
+        produces `v_first` rather than mixing towards it and never reads that chain.
+        """
+        w_log = -_INV_SQRT_E * torch.sigmoid(
+            _rwkv7_matmul(torch.tanh(_rwkv7_matmul(xw, self.w1, tokenwise)), self.w2, tokenwise) + self.w0
+        )
+        a = torch.sigmoid(_rwkv7_matmul(_rwkv7_matmul(xa, self.a1, tokenwise), self.a2, tokenwise) + self.a0)
+        g = _rwkv7_matmul(torch.sigmoid(_rwkv7_matmul(xg, self.g1, tokenwise)), self.g2, tokenwise)
+        v_gate = (
+            None
+            if xv is None
+            else torch.sigmoid(_rwkv7_matmul(_rwkv7_matmul(xv, self.v1, tokenwise), self.v2, tokenwise) + self.v0)
+        )
+        return w_log, a, g, v_gate
+
+
+class Rwkv7FeedForward(nn.Module):
+    """RWKV-7 channel-mixing block: squared-ReLU over a single token shift."""
+
+    def __init__(self, config: Rwkv7Config, layer_id: int):
+        super().__init__()
+        self.config = config
+        self.layer_id = layer_id
+        C = config.hidden_size
+        self.time_shift = Rwkv7TokenShift()
+        self.x_k = nn.Parameter(torch.zeros(1, 1, C))
+        self.key = nn.Linear(C, config.intermediate_size, bias=False)
+        self.value = nn.Linear(config.intermediate_size, C, bias=False)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        shift_state: torch.Tensor | None,
+        cu_seq_lens: torch.Tensor | None = None,
+        keep: torch.Tensor | None = None,
+    ):
+        tokenwise = self.config.wkv_implementation != "chunked"
+        shifted, new_shift_state = self.time_shift(hidden_states, shift_state, cu_seq_lens, keep)
+        xk = hidden_states + self.x_k * (shifted - hidden_states)
+        inner = torch.relu(_rwkv7_linear(self.key, xk, tokenwise)) ** 2
+        out = self._project(inner, tokenwise)
+        return out, new_shift_state
+
+    def _project(self, inner: torch.Tensor, tokenwise: bool) -> torch.Tensor:
+        return _rwkv7_linear(self.value, inner, tokenwise)
+
+
+class Rwkv7CacheLayer(LinearAttentionLayer):
+    """One block's slice of the recurrent state: the WKV matrix and two token shifts.
+
+    Everything a beam search or a batched generate needs to do to a cache is a
+    permutation of its batch axis, and for RWKV-7 that is the whole job -- the state
+    is O(1) in sequence length, so there is no time axis to gather along and no
+    length bookkeeping to keep consistent.
+
+    The slot layout is [`Rwkv7Cache`]'s; this class only moves whatever is in them.
+    """
+
+    def lazy_initialization(self, conv_states=None, recurrent_states=None, state_idx: int = 0, **kwargs) -> None:
+        super().lazy_initialization(conv_states, recurrent_states, state_idx, **kwargs)
+        # Upstream records device/dtype only on the conv branch, since a linear
+        # attention layer normally has a convolution in front of it. This model has
+        # none, so without this every recurrent-only layer would keep `device=None`
+        # and `reorder_cache`'s `beam_idx.to(self.device)` would fail.
+        if recurrent_states is not None and self.device is None:
+            self.dtype, self.device = recurrent_states.dtype, recurrent_states.device
+
+    def allocate(
+        self, batch: int, shapes: dict[int, tuple], device, dtypes: dict[int, torch.dtype], state_dtype
+    ) -> None:
+        """Create every slot up front, zeroed, at pinned addresses.
+
+        The lazy path allocates a slot the first time it is written, which is inside
+        the compiled region -- and a buffer that first appears there cannot be given
+        a static address, so inductor declines CUDA graphs for a recurrent decode,
+        the one workload that most needs them.
+        """
+        for slot, shape in shapes.items():
+            buffer = torch.zeros((batch, *shape), device=device, dtype=dtypes[slot])
+            if not torch.compiler.is_compiling():
+                torch._dynamo.mark_static_address(buffer)
+            self.recurrent_states[slot] = buffer
+            self.is_recurrent_states_initialized[slot] = True
+            self.has_previous_state[slot] = True
+        self.dtype, self.device = state_dtype, device
+
+    def reorder_cache(self, beam_idx: torch.LongTensor) -> None:
+        """Permute the batch onto the beams that survived, without moving house.
+
+        Upstream rebinds each slot to the `index_select` result. That is correct but
+        it hands back a freshly allocated tensor, so the address pinned at allocation
+        is gone for the rest of the generation and the compiled decode quietly loses
+        its CUDA graphs. Copying back into the same buffer costs one temporary and
+        keeps the pinning.
+        """
+        for slot in range(self.number_of_states):
+            buffer = self.recurrent_states[slot]
+            if self.is_recurrent_states_initialized[slot] and buffer is not None:
+                buffer.copy_(buffer.index_select(0, beam_idx.to(buffer.device)))
+
+    def batch_repeat_interleave(self, repeats: int) -> None:
+        """Fan each sequence's state out to `repeats` copies, for one-prompt-many-samples.
+
+        This changes the batch size, so unlike `reorder_cache` it cannot preserve the
+        pinned addresses -- the buffers are necessarily new ones. Callers that then
+        compile should re-pin, which `Rwkv7Model.allocate_state` does.
+        """
+        for slot in range(self.number_of_states):
+            buffer = self.recurrent_states[slot]
+            if self.is_recurrent_states_initialized[slot] and buffer is not None:
+                self.recurrent_states[slot] = buffer.repeat_interleave(repeats, dim=0)
+
+    def batch_select_indices(self, indices: torch.Tensor) -> None:
+        """Keep only the given batch rows. Same address caveat as `batch_repeat_interleave`."""
+        for slot in range(self.number_of_states):
+            buffer = self.recurrent_states[slot]
+            if self.is_recurrent_states_initialized[slot] and buffer is not None:
+                self.recurrent_states[slot] = buffer[indices, ...]
+
+
+class Rwkv7Cache(Cache):
+    """The recurrent state of every block, as a `Cache`.
+
+    It replaces a KV cache and is a constant size: one `[num_heads, head_dim,
+    head_dim]` matrix and two `[hidden]` token shifts per layer per sequence,
+    whatever the context length. That is the property the architecture is for, so
+    `get_max_length()` is -1 (no limit) and nothing here grows as tokens arrive.
+    """
+
+    # Slot numbering inside one block's layer. All three are recurrent states:
+    # RWKV-7's token shift is a one-token history, not a convolution window, so it
+    # lives in a recurrent slot, where the update is a plain copy rather than the
+    # rolling concatenate a conv slot would do.
+    WKV, ATT_SHIFT, FFN_SHIFT = 0, 1, 2
+
+    def __init__(
+        self,
+        config: Rwkv7Config,
+        batch_size: int | None = None,
+        device: torch.device | str | None = None,
+        dtype: torch.dtype | None = None,
+    ):
+        super().__init__(layers=[Rwkv7CacheLayer(number_of_states=3) for _ in range(config.num_hidden_layers)])
+        self.config = config
+        if batch_size is not None:
+            self.allocate(batch_size, device, dtype)
+
+    def allocate(self, batch_size: int, device=None, dtype=None) -> "Rwkv7Cache":
+        config = self.config
+        shapes = {
+            self.WKV: (config.num_heads, config.head_dim, config.head_dim),
+            self.ATT_SHIFT: (config.hidden_size,),
+            self.FFN_SHIFT: (config.hidden_size,),
+        }
+        # The WKV state carries the whole history and is the one place where a
+        # narrower dtype actually costs accuracy, so it is configured separately
+        # from the activation dtype the shifts follow.
+        dtypes = {
+            self.WKV: getattr(torch, config.wkv_state_dtype),
+            self.ATT_SHIFT: dtype,
+            self.FFN_SHIFT: dtype,
+        }
+        for layer in self.layers:
+            assert isinstance(layer, Rwkv7CacheLayer)
+            layer.allocate(batch_size, shapes, device, dtypes, dtypes[self.WKV])
+        return self
+
+    def read(self, layer_idx: int):
+        """This block's `(att_shift, ffn_shift, wkv)`, each None before allocation."""
+        layer = self.layers[layer_idx]
+        assert isinstance(layer, Rwkv7CacheLayer)
+        states = layer.recurrent_states
+        return states[self.ATT_SHIFT], states[self.FFN_SHIFT], states[self.WKV]
+
+    def write(self, layer_idx: int, att_shift: torch.Tensor, ffn_shift: torch.Tensor, wkv: torch.Tensor) -> None:
+        states = (
+            (self.ATT_SHIFT, att_shift),
+            (self.FFN_SHIFT, ffn_shift),
+            (self.WKV, wkv),
+        )
+        # Copying into the pre-allocated slot is what keeps its address fixed, which
+        # is what lets a captured CUDA graph replay the decode loop. It is also an
+        # in-place write on a tensor autograd may be holding, and backward through a
+        # cached forward then dies on "a variable needed for gradient computation has
+        # been modified by an inplace operation". Rebinding instead costs nothing
+        # here, because a training step is not the workload that wants graphs.
+        if torch.is_grad_enabled() and any(state.requires_grad for _, state in states):
+            layer = self.layers[layer_idx]
+            assert isinstance(layer, Rwkv7CacheLayer)
+            for slot, state in states:
+                layer.recurrent_states[slot] = state
+                layer.is_recurrent_states_initialized[slot] = True
+                layer.has_previous_state[slot] = True
+            return
+        for slot, state in states:
+            self.update_recurrent_state(state, layer_idx, slot)
+
+
+class Rwkv7Block(GradientCheckpointingLayer):
+    def __init__(self, config: Rwkv7Config, layer_id: int):
+        super().__init__()
+        self.config = config
+        self.layer_id = layer_id
+        eps, bias = config.norm_eps, config.norm_bias
+        # Layer 0 carries the extra input norm of the reference implementation.
+        if layer_id == 0:
+            self.ln0 = nn.LayerNorm(config.hidden_size, eps=eps, bias=bias)
+        self.ln1 = nn.LayerNorm(config.hidden_size, eps=eps, bias=bias)
+        self.ln2 = nn.LayerNorm(config.hidden_size, eps=eps, bias=bias)
+        self.att = Rwkv7Attention(config, layer_id)
+        self.ffn = Rwkv7FeedForward(config, layer_id)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        v_first: torch.Tensor | None,
+        state: Rwkv7Cache | None,
+        keep: torch.Tensor | None = None,
+        cu_seq_lens: torch.Tensor | None = None,
+    ):
+        if self.layer_id == 0:
+            hidden_states = self.ln0(hidden_states)
+
+        att_shift, ffn_shift, wkv = state.read(self.layer_id) if state is not None else (None, None, None)
+
+        # Padding is blanked AFTER each norm, not before: a LayerNorm maps the zero
+        # vector to its own bias, so masking the residual stream instead would let
+        # every pad position come back to life on the way into the next mixer -- and
+        # a live pad both moves the state and leaks into the next token's shift.
+        attn_in = self.ln1(hidden_states)
+        if keep is not None:
+            attn_in = attn_in * keep
+        attn_out, v_first, att_shift, wkv = self.att(attn_in, v_first, att_shift, wkv, keep, cu_seq_lens)
+        hidden_states = hidden_states + attn_out
+
+        ffn_in = self.ln2(hidden_states)
+        if keep is not None:
+            ffn_in = ffn_in * keep
+        ffn_out, ffn_shift = self.ffn(ffn_in, ffn_shift, cu_seq_lens, keep)
+        hidden_states = hidden_states + ffn_out
+
+        if state is not None:
+            # `update_recurrent_state` copies into the pre-allocated slot rather than
+            # rebinding it, so the buffers keep fixed addresses across steps -- which
+            # is what lets a captured CUDA graph replay the decode loop.
+            state.write(self.layer_id, att_shift, ffn_shift, wkv)
+        return hidden_states, v_first, state
+
+
+@dataclass
+class Rwkv7Output(ModelOutput):
+    r"""
+    state (`Rwkv7Cache`, *optional*):
+        The recurrent state of every block. Feed it back to continue a sequence; it
+        replaces the KV cache and is a constant size, whatever the context length.
+    attentions (always `None`):
+        Present so that code written against the common output shape does not have to
+        special-case this model. There is no attention here to report -- the mixing is
+        a recurrence, not a score matrix -- so it is `None` whatever `output_attentions`
+        is set to, rather than an empty tuple pretending to be a per-layer list.
+    """
+
+    last_hidden_state: torch.FloatTensor | None = None
+    state: Rwkv7Cache | None = None
+    hidden_states: tuple[torch.FloatTensor, ...] | None = None
+    attentions: None = None
+
+
+@dataclass
+class Rwkv7CausalLMOutput(ModelOutput):
+    r"""
+    state (`Rwkv7Cache`, *optional*):
+        The recurrent state, as in [`Rwkv7Output`].
+    """
+
+    loss: torch.FloatTensor | None = None
+    logits: torch.FloatTensor | None = None
+    state: Rwkv7Cache | None = None
+    hidden_states: tuple[torch.FloatTensor, ...] | None = None
+    attentions: None = None
+
+
+@auto_docstring
+class Rwkv7PreTrainedModel(PreTrainedModel):
+    config: Rwkv7Config
+    base_model_prefix = "rwkv7"
+    _no_split_modules = ["Rwkv7Block"]
+    supports_gradient_checkpointing = True
+    _is_stateful = True
+    # Beam search reorders through `Rwkv7Cache.reorder_cache`. Defining
+    # `_reorder_cache` here instead would take precedence over it in `generate` and
+    # bypass the cache's own bookkeeping, so it is deliberately absent.
+
+    @torch.no_grad()
+    def _init_weights(self, module):
+        """Reproduce RWKV-LM's x070 constructor and final `generate_init_weight` pass."""
+        super()._init_weights(module)
+        if isinstance(module, Rwkv7Attention):
+            layer_id = module.layer_id
+            num_layers = module.config.num_hidden_layers
+            hidden_size = module.config.hidden_size
+            head_dim = module.config.head_dim
+
+            ratio_0_to_1 = layer_id / (num_layers - 1) if num_layers > 1 else 0.0
+            ratio_1_to_almost0 = 1.0 - layer_id / num_layers
+            channel = torch.arange(hidden_size, dtype=torch.float32, device=module.x_r.device)
+            ddd = (channel / hidden_size).reshape(1, 1, hidden_size)
+            linear = channel / max(hidden_size - 1, 1) - 0.5
+            if head_dim > 1:
+                zigzag = ((channel % head_dim) - (head_dim - 1) / 2) / ((head_dim - 1) / 2)
+                zigzag = zigzag * zigzag.abs()
+            else:
+                zigzag = torch.zeros_like(channel)
+            time_decay = -6 + 6 * (channel / max(hidden_size - 1, 1)) ** (1 + ratio_0_to_1**0.3)
+
+            init.copy_(module.x_r, 1 - ddd ** (0.2 * ratio_1_to_almost0))
+            init.copy_(module.x_w, 1 - ddd ** (0.9 * ratio_1_to_almost0))
+            init.copy_(module.x_k, 1 - ddd ** (0.7 * ratio_1_to_almost0))
+            init.copy_(module.x_v, 1 - ddd ** (0.7 * ratio_1_to_almost0))
+            init.copy_(module.x_a, 1 - ddd ** (0.9 * ratio_1_to_almost0))
+            init.copy_(module.x_g, 1 - ddd ** (0.2 * ratio_1_to_almost0))
+
+            init.zeros_(module.w1)
+            _rwkv7_lora_orthogonal_(module.w2)
+            init.copy_(module.w0, (time_decay + 0.5 + zigzag * 2.5).reshape(1, 1, hidden_size))
+            init.zeros_(module.a1)
+            _rwkv7_lora_orthogonal_(module.a2)
+            init.copy_(module.a0, (-0.19 + zigzag * 0.3 + linear * 0.4).reshape(1, 1, hidden_size))
+            init.zeros_(module.v1)
+            _rwkv7_lora_orthogonal_(module.v2)
+            init.copy_(module.v0, (0.73 - linear * 0.4).reshape(1, 1, hidden_size))
+            init.zeros_(module.g1)
+            _rwkv7_lora_orthogonal_(module.g2)
+            init.copy_(module.k_k, (0.71 - linear * 0.1).reshape(1, 1, hidden_size))
+            init.constant_(module.k_a, 1.02)
+            init.constant_(module.r_k, -0.04)
+
+            init.orthogonal_(module.key.weight, gain=0.1)
+            init.zeros_(module.output.weight)
+            init.constant_(module.ln_x.weight, ((layer_id + 1) / num_layers) ** 0.7)
+            if module.ln_x.bias is not None:
+                init.zeros_(module.ln_x.bias)
+        elif isinstance(module, Rwkv7FeedForward):
+            ratio_1_to_almost0 = 1.0 - module.layer_id / module.config.num_hidden_layers
+            channel = torch.arange(module.config.hidden_size, dtype=torch.float32, device=module.x_k.device)
+            ddd = (channel / module.config.hidden_size).reshape(1, 1, module.config.hidden_size)
+            init.copy_(module.x_k, 1 - ddd ** (ratio_1_to_almost0**4))
+            init.zeros_(module.value.weight)
+        elif isinstance(module, nn.Linear):
+            init.orthogonal_(module.weight)
+            if module.bias is not None:
+                init.zeros_(module.bias)
+        elif isinstance(module, nn.Embedding):
+            init.uniform_(module.weight, a=-1e-4, b=1e-4)
+
+
+@auto_docstring
+class Rwkv7Model(Rwkv7PreTrainedModel):
+    def __init__(self, config: Rwkv7Config):
+        super().__init__(config)
+        self.emb = nn.Embedding(config.vocab_size, config.hidden_size)
+        self.blocks = nn.ModuleList([Rwkv7Block(config, i) for i in range(config.num_hidden_layers)])
+        self.ln_out = nn.LayerNorm(config.hidden_size, eps=config.norm_eps, bias=config.norm_bias)
+        self.gradient_checkpointing = False
+        self.post_init()
+
+    def get_input_embeddings(self):
+        return self.emb
+
+    def set_input_embeddings(self, new_embeddings):
+        self.emb = new_embeddings
+
+    def allocate_state(self, batch: int, device=None, dtype=None) -> Rwkv7Cache:
+        """A zeroed cache, plus everything else the decode path would build lazily.
+
+        Call this before compiling, and pass the result in as `state=`.
+        `mark_static_address` cannot run during tracing, so anything first allocated
+        *inside* the compiled region stays unpinned, and inductor declines CUDA graphs
+        for a region that mutates its inputs. Starting from `state=None` loses the
+        state buffers that way, and the decode then runs several times slower while
+        saying so only in a line of warning.
+        """
+        state = self._empty_state(
+            batch,
+            device if device is not None else self.emb.weight.device,
+            dtype if dtype is not None else self.emb.weight.dtype,
+        )
+        return state
+
+    def _empty_state(self, batch: int, device, dtype) -> Rwkv7Cache:
+        return Rwkv7Cache(self.config, batch_size=batch, device=device, dtype=dtype)
+
+    @can_return_tuple
+    @auto_docstring
+    def forward(
+        self,
+        input_ids: torch.LongTensor | None = None,
+        attention_mask: torch.LongTensor | None = None,
+        inputs_embeds: torch.FloatTensor | None = None,
+        state: Rwkv7Cache | None = None,
+        cu_seq_lens: torch.LongTensor | None = None,
+        use_cache: bool | None = None,
+        output_attentions: bool | None = None,
+        output_hidden_states: bool | None = None,
+        **kwargs,
+    ) -> Rwkv7Output:
+        r"""
+        attention_mask (`torch.LongTensor`, *optional*):
+            1 on real tokens, 0 on padding. Read as the **tail** of whatever is
+            given, so a decode step may hand over the whole conversation's mask and
+            only the last position is used. A prefix chunk must therefore slice its
+            own mask to match the `input_ids` it passes, or it silently masks the
+            wrong positions. There is nothing here for a mask to hide behind: an
+            all-recurrent model feeds pads through the recurrence like any other
+            token unless they are neutralised.
+        state (`Rwkv7Cache`, *optional*):
+            Recurrent state returned by a previous call; pass it back to continue
+            the sequence. Allocated on the first forward if omitted. Ignored as a
+            *history* when `cu_seq_lens` is given -- see there.
+        cu_seq_lens (`torch.LongTensor`, *optional*):
+            Cumulative sequence lengths for a *packed* batch: several sequences
+            concatenated into one row instead of padded to a rectangle, starting at
+            0 and ending at `seq_len`, and non-decreasing. Each segment then decodes
+            from a fresh recurrent state, as if it had been run on its own. This is
+            the varlen layout; use it instead of padding when the lengths vary a
+            lot, since a recurrent model pays for pad tokens in time as well as
+            memory.
+
+            A packed batch is a set of *new* sequences, not a continuation, so a
+            `state` passed alongside contributes its shape and dtype and nothing
+            else -- its contents are not read. What comes back is the last segment's
+            state, which is the one a continuation of this row would resume from.
+        """
+        output_hidden_states = (
+            output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
+        )
+        use_cache = use_cache if use_cache is not None else self.config.use_cache
+        # Gradient checkpointing replays the forward during backward. If the cache is
+        # live, the replay re-reads a state the first pass already advanced, and the
+        # gradients that come back are wrong rather than absent. `GradientCheckpointingLayer`
+        # neutralises this itself, but only for a cache arriving as a keyword named
+        # `use_cache` / `past_key_values` / `layer_past`; this model hands its state to
+        # the block positionally, so none of those guards see it and the whole thing
+        # has to be caught here.
+        if self.gradient_checkpointing and self.training and use_cache:
+            logger.warning_once(
+                "`use_cache=True` is incompatible with gradient checkpointing. Setting `use_cache=False`."
+            )
+            use_cache = False
+
+        if (input_ids is None) == (inputs_embeds is None):
+            raise ValueError("Specify exactly one of input_ids or inputs_embeds")
+        if inputs_embeds is None:
+            inputs_embeds = self.emb(input_ids)
+
+        if use_cache and state is None:
+            state = self._empty_state(inputs_embeds.shape[0], inputs_embeds.device, inputs_embeds.dtype)
+
+        # An attention-free recurrence has to be told where the padding is: pads are
+        # fed through the recurrence like any other token, so a left-padded batch
+        # (what `generate` produces) would otherwise start every short row from a
+        # state the pads had already moved. A single decoded token is never padding,
+        # so the decode step takes none of this.
+        keep = None
+        if attention_mask is not None:
+            # The mask is taken as the TAIL of whatever is passed, so a decode step
+            # can hand over the whole conversation's mask. A prefix chunk must slice
+            # its own mask to match, or it silently reads the wrong positions.
+            #
+            # Applied at seq_len == 1 as well. Skipping it there assumes a single
+            # decoded token is never padding -- true of `generate`, but an
+            # assumption about the caller rather than a property
+            # of the model, and a fully-masked 1-token row was moving the state.
+            keep = attention_mask[:, -inputs_embeds.shape[1] :, None].to(inputs_embeds.dtype)
+
+        if cu_seq_lens is not None:
+            # Checked rather than trusted: a malformed boundary list does not fail,
+            # it silently splits the recurrence in the wrong places and returns
+            # fluent output computed from the wrong states.
+            if inputs_embeds.shape[0] != 1:
+                raise ValueError(
+                    f"cu_seq_lens describes one packed row, but got batch size {inputs_embeds.shape[0]}. "
+                    "Pack the sequences into a single row, or use attention_mask with a padded batch."
+                )
+            if cu_seq_lens.ndim != 1 or cu_seq_lens[0] != 0 or cu_seq_lens[-1] != inputs_embeds.shape[1]:
+                raise ValueError(
+                    f"cu_seq_lens must be 1-D, start at 0 and end at seq_len ({inputs_embeds.shape[1]}); "
+                    f"got {cu_seq_lens.tolist()}"
+                )
+            # Endpoints alone do not pin the list down. A pair that goes backwards
+            # is skipped rather than rejected further in, so the segments emit fewer
+            # tokens than came in and the row silently changes length -- and one that
+            # merely repeats a boundary contributes an empty segment, which is
+            # harmless but is never what the caller meant.
+            if bool((cu_seq_lens[1:] <= cu_seq_lens[:-1]).any()):
+                raise ValueError(
+                    "cu_seq_lens must be strictly increasing (each segment needs at least one token); "
+                    f"got {cu_seq_lens.tolist()}"
+                )
+
+        hidden_states = inputs_embeds
+        v_first = None
+        all_hidden_states = () if output_hidden_states else None
+
+        for block in self.blocks:
+            if output_hidden_states:
+                all_hidden_states = all_hidden_states + (hidden_states,)
+            hidden_states, v_first, state = block(hidden_states, v_first, state, keep, cu_seq_lens)
+
+        hidden_states = self.ln_out(hidden_states)
+        if output_hidden_states:
+            all_hidden_states = all_hidden_states + (hidden_states,)
+
+        return Rwkv7Output(last_hidden_state=hidden_states, state=state, hidden_states=all_hidden_states)
+
+
+@auto_docstring
+class Rwkv7ForCausalLM(Rwkv7PreTrainedModel, GenerationMixin):
+    _tied_weights_keys = {"head.weight": "rwkv7.emb.weight"}
+
+    def __init__(self, config: Rwkv7Config):
+        super().__init__(config)
+        self.rwkv7 = Rwkv7Model(config)
+        self.head = nn.Linear(config.hidden_size, config.vocab_size, bias=False)
+        self.post_init()
+
+    @torch.no_grad()
+    def _init_weights(self, module):
+        super()._init_weights(module)
+        # Composite-model dispatch sends the base RWKV7 modules through their own
+        # initializer, so every Linear reaching this override is an LM head. Using
+        # its actual shape also covers a newly resized, not-yet-installed head.
+        if isinstance(module, nn.Linear):
+            gain = 0.5 * math.sqrt(module.out_features / module.in_features)
+            if module.out_features <= module.in_features:
+                gain = 0.5
+            init.orthogonal_(module.weight, gain=gain)
+
+    def get_output_embeddings(self):
+        return self.head
+
+    def set_output_embeddings(self, new_embeddings):
+        self.head = new_embeddings
+
+    def prepare_inputs_for_generation(
+        self,
+        input_ids,
+        state=None,
+        inputs_embeds=None,
+        next_sequence_length=None,
+        is_first_iteration=False,
+        **kwargs,
+    ):
+        # `state is not None` does not by itself mean decode: callers can provide
+        # an empty preallocated state for the initial prompt, or a carried state
+        # followed by a multi-token continuation. GenerationMixin tells us how
+        # much of the input is new through `next_sequence_length`.
+        if next_sequence_length is not None:
+            input_ids = input_ids[:, -next_sequence_length:]
+            if inputs_embeds is not None:
+                inputs_embeds = inputs_embeds[:, -next_sequence_length:]
+        elif state is not None and not is_first_iteration:
+            input_ids = input_ids[:, -1:]
+        model_inputs = (
+            {"input_ids": input_ids}
+            if inputs_embeds is None or not is_first_iteration
+            else {"inputs_embeds": inputs_embeds}
+        )
+        model_inputs["state"] = state
+        # Everything else the caller passed goes through, minus what this model does
+        # not take. An allowlist of specific names instead would make
+        # `generate(output_hidden_states=True)` return a tuple of `None`: the flag
+        # would be dropped here, the forward would never see it, and generate collects the
+        # nothing it got back. Any user kwarg met the same fate, silently. The two
+        # excluded here are `labels`, which would make generate compute a loss it
+        # never reads, and the KV-cache bookkeeping that belongs to models with a KV
+        # cache; this one carries its history in `state`.
+        skip = ("labels", "past_key_values", "cache_position")
+        model_inputs.update({k: v for k, v in kwargs.items() if k not in skip and k not in model_inputs})
+        return model_inputs
+
+    @can_return_tuple
+    @auto_docstring
+    def forward(
+        self,
+        input_ids: torch.LongTensor | None = None,
+        attention_mask: torch.LongTensor | None = None,
+        inputs_embeds: torch.FloatTensor | None = None,
+        state: Rwkv7Cache | None = None,
+        labels: torch.LongTensor | None = None,
+        use_cache: bool | None = None,
+        output_attentions: bool | None = None,
+        output_hidden_states: bool | None = None,
+        logits_to_keep: int | torch.Tensor = 0,
+        **kwargs,
+    ) -> Rwkv7CausalLMOutput:
+        r"""
+        state (`Rwkv7Cache`, *optional*):
+            Recurrent state returned by a previous call.
+        logits_to_keep (`int` or `torch.Tensor`, *optional*, defaults to 0):
+            Compute the head over only the last `logits_to_keep` positions, or over the
+            positions this tensor indexes; 0 means all of them. Worth having on a
+            recurrent model for the same reason as on any other: a prefill needs one
+            row of logits and the vocabulary is the widest matrix in the model, so
+            running the head over the whole prompt is the single largest avoidable cost
+            in a prefill. Until this argument existed it was swallowed by `**kwargs`,
+            so `generate` declined to pass it and a caller who passed it was quietly
+            ignored.
+        """
+        outputs = self.rwkv7(
+            input_ids=input_ids,
+            attention_mask=attention_mask,
+            inputs_embeds=inputs_embeds,
+            state=state,
+            use_cache=use_cache,
+            output_hidden_states=output_hidden_states,
+            **kwargs,
+        )
+        keep = slice(-logits_to_keep, None) if isinstance(logits_to_keep, int) else logits_to_keep
+        hidden_to_project = outputs.last_hidden_state[:, keep, :]
+        logits = _rwkv7_linear(
+            self.head,
+            hidden_to_project,
+            self.config.wkv_implementation != "chunked",
+        )
+
+        loss = None
+        if labels is not None:
+            loss = self.loss_function(logits, labels, self.config.vocab_size, **kwargs)
+
+        return Rwkv7CausalLMOutput(loss=loss, logits=logits, state=outputs.state, hidden_states=outputs.hidden_states)
+
+
+__all__ = ["Rwkv7Cache", "Rwkv7PreTrainedModel", "Rwkv7Model", "Rwkv7ForCausalLM"]

release-manifest.json

@@ -0,0 +1,332 @@
+{
+  "builder": {
+    "asset_set": "2026.08.06-r1",
+    "assets_sha256": "3c3f6237fcd9dc2ec8ba0b61b626ef72924ebe42bdc6dc49e21a5bfeb19bb96b",
+    "version": "0.2.0"
+  },
+  "conversion": {
+    "explicit_cast": false,
+    "serialized_parameter_count": 1527668736,
+    "source_float_dtypes": [
+      "bfloat16"
+    ],
+    "source_parameter_count": 1527668736,
+    "synthesized_tensors": [],
+    "target_float_dtypes": [
+      "bfloat16"
+    ],
+    "tensor_count": 798,
+    "tensor_map_sha256": "03131ced241e7b0f86869363b3969ead6ef58462f23d9f04e56e26d00565aefc"
+  },
+  "files": {
+    ".gitattributes": {
+      "role": "metadata",
+      "sha256": "edfb97afcbabd3c89eb1d19070521d403f4644ecacfec826e14c2dc68cf419d8",
+      "size_bytes": 134
+    },
+    "LICENSE": {
+      "role": "metadata",
+      "sha256": "475641f33cfd755f19245d2eb6a6d8fb7fdca60835cfa369b7e3774460409cc4",
+      "size_bytes": 11357
+    },
+    "NOTICE": {
+      "role": "metadata",
+      "sha256": "7f8e7f0f3bf0a8e7caf13132770cfe841e5eec10203c18798f8051e96d234108",
+      "size_bytes": 250
+    },
+    "README.md": {
+      "role": "model_card",
+      "sha256": "b7795213e1fcf1d9e7de41fa027a0b2d311b2dacfa5a5ba9f62c4c2484bd729f",
+      "size_bytes": 10260
+    },
+    "chat_template.jinja": {
+      "role": "tokenizer",
+      "sha256": "f0d8f8161f03d38d0bf68b1ea9a1293064c0fe2b60dfc2874008154c33861c07",
+      "size_bytes": 2333
+    },
+    "config.json": {
+      "role": "model_config",
+      "sha256": "442d1155144fd8331bac86250131cc1a874d8d41c1504f930e0b522a99b615aa",
+      "size_bytes": 748
+    },
+    "configuration_rwkv7.py": {
+      "role": "model_code",
+      "sha256": "a4a2adcd02a101cc2a17f2452f8e56f4a8d8348be0d19ab989e908a83921c123",
+      "size_bytes": 7318
+    },
+    "generation_config.json": {
+      "role": "model_config",
+      "sha256": "358114ae91c5eee6fe10b41a22a9ee95a9b1c20ed86d0666617aa049386d63a3",
+      "size_bytes": 87
+    },
+    "inference/generate.py": {
+      "role": "inference",
+      "sha256": "9e4f3461ade971a96a276ec4e18683d6d8df6420b09bf3e63cc499b6e119aa05",
+      "size_bytes": 7804
+    },
+    "inference/kernel.py": {
+      "role": "inference",
+      "sha256": "88379d8979b908a01af1f06b89d8fc624d332796db6044ec4a027d7cecab1cb6",
+      "size_bytes": 165253
+    },
+    "inference/model_loader.py": {
+      "role": "inference",
+      "sha256": "c48abe28abde91ede570498372cd6b9aa6a05080525df7f5045613849ed9bb61",
+      "size_bytes": 6452
+    },
+    "inference/requirements.txt": {
+      "role": "inference",
+      "sha256": "9cc4f89e2edc45da7a04cdf8e4c1da83b682a6bf969e0eed7b3dec91128d925a",
+      "size_bytes": 93
+    },
+    "inference/runtime.py": {
+      "role": "inference",
+      "sha256": "064c53ce5dfe811ea873ad9976f8f0839b55e515e604adf3c3e28fd4c02de784",
+      "size_bytes": 158191
+    },
+    "model.safetensors": {
+      "role": "weights",
+      "sha256": "84ccbb857c84e00cefc48b233937ada79c411e491df25fb21aed23237f39a14f",
+      "size_bytes": 3055418240
+    },
+    "modeling_rwkv7.py": {
+      "role": "model_code",
+      "sha256": "f3a1d9bde755f54887451e8755c6d48866ecc9146610c7d8fa73e0fa68c6734c",
+      "size_bytes": 57446
+    },
+    "tokenizer.json": {
+      "role": "tokenizer",
+      "sha256": "1c0bc8e5f355de60a4587a1e9abb9aa4ef65508ff59641344e7f9f10f76feacf",
+      "size_bytes": 1580973
+    },
+    "tokenizer_config.json": {
+      "role": "tokenizer",
+      "sha256": "58a994975a1a64327cd7ee71f230d4a5f290e7260957b18840e9ecf0d5d57d0e",
+      "size_bytes": 407
+    }
+  },
+  "identity": {
+    "parameter_label": "1.5",
+    "release_date": "20260805",
+    "training_context_length": 16384
+  },
+  "metadata": {
+    "context_length": 16384,
+    "datasets": [
+      "HuggingFaceFW/fineweb-edu",
+      "mlfoundations/dclm-baseline-1.0",
+      "cerebras/SlimPajama-627B",
+      "EleutherAI/pile",
+      "bigcode/starcoderdata",
+      "oscar-corpus/OSCAR-2301"
+    ],
+    "languages": [
+      "en",
+      "zh",
+      "fr",
+      "es",
+      "de",
+      "pt",
+      "ru",
+      "it",
+      "ja",
+      "ko",
+      "vi",
+      "ar"
+    ],
+    "license": "apache-2.0",
+    "profile": "g1i",
+    "provenance": "locked-profile"
+  },
+  "model_code": {
+    "format_version": 1,
+    "source_repository": "https://github.com/huggingface/transformers.git",
+    "source_revision": "4ad9ed0747ed6ba75c787e8f9040dcd64b166ee2",
+    "sources": {
+      "configuration_rwkv7.py": {
+        "asset_path": "model_code/configuration_rwkv7.py",
+        "output_sha256": "a4a2adcd02a101cc2a17f2452f8e56f4a8d8348be0d19ab989e908a83921c123",
+        "repository_path": "src/transformers/models/rwkv7/configuration_rwkv7.py",
+        "source_sha256": "6f5b92c5fe7498ad22b0054a2f735a7ca82e7577436f4ad32f0fc27d1e900fdd"
+      },
+      "modeling_rwkv7.py": {
+        "asset_path": "model_code/modeling_rwkv7.py",
+        "output_sha256": "f3a1d9bde755f54887451e8755c6d48866ecc9146610c7d8fa73e0fa68c6734c",
+        "repository_path": "src/transformers/models/rwkv7/modeling_rwkv7.py",
+        "source_sha256": "3e8e5af7c4eba0b5de1496aef44773d7ac1bb4d96756e6f55efaf29453d67952"
+      }
+    },
+    "transformers_min_version": "5.15"
+  },
+  "profile": {
+    "checkpoint": "g1i-1.5b-20260805",
+    "family": "g1i"
+  },
+  "runtime": {
+    "files": {
+      "kernel.py": {
+        "sha256": "88379d8979b908a01af1f06b89d8fc624d332796db6044ec4a027d7cecab1cb6"
+      },
+      "runtime.py": {
+        "sha256": "064c53ce5dfe811ea873ad9976f8f0839b55e515e604adf3c3e28fd4c02de784"
+      }
+    },
+    "format_version": 4,
+    "kernel": {
+      "format_version": 3,
+      "output": "inference/kernel.py",
+      "output_sha256": "88379d8979b908a01af1f06b89d8fc624d332796db6044ec4a027d7cecab1cb6",
+      "sources": {
+        "kernel_tilelang_decode.py": {
+          "exports": [
+            "Any",
+            "_HEAD_SIZE",
+            "_NEXP_HALF_LOG2_E",
+            "_NLOG2_E",
+            "_ROTATOR1_SIGNED",
+            "_TWO_NEG_41",
+            "_build_cmix_add_layernorm_mix_program",
+            "_build_cmix_binned_finalize_program",
+            "_build_cmix_finalize_program",
+            "_build_cmix_layernorm_mix_program",
+            "_build_cmix_sparse_atomic_program",
+            "_build_cmix_sparse_binned_program",
+            "_build_cmix_sparse_split_program",
+            "_build_cmix_value_program",
+            "_build_ffn_program",
+            "_build_gemv_program",
+            "_build_key_gate_program",
+            "_build_post_state_program",
+            "_build_rankout_program",
+            "_build_rankout_reduced_program",
+            "_build_rkv_program",
+            "_build_tmix_layernorm_mix6_program",
+            "_build_wkv_program",
+            "_build_wkv_w0_t1_program",
+            "_cmix_value_out",
+            "_compiled_cmix_add_layernorm_mix",
+            "_compiled_cmix_binned_finalize",
+            "_compiled_cmix_finalize",
+            "_compiled_cmix_layernorm_mix",
+            "_compiled_cmix_sparse_atomic",
+            "_compiled_cmix_sparse_binned",
+            "_compiled_cmix_sparse_split",
+            "_compiled_cmix_value",
+            "_compiled_ffn",
+            "_compiled_gemv",
+            "_compiled_key_gate",
+            "_compiled_post_state",
+            "_compiled_rankout",
+            "_compiled_rankout_reduced",
+            "_compiled_rkv",
+            "_compiled_tmix_layernorm_mix6",
+            "_compiled_wkv",
+            "_compiled_wkv_w0_t1",
+            "_dtype_name",
+            "_ffn_out",
+            "_post_state_out",
+            "_rankout_out",
+            "_require_contiguous_cuda",
+            "_rkv_out",
+            "_wkv_kernel_out",
+            "_wkv_out",
+            "_wkv_precise_out",
+            "_wkv_w0_t1_out",
+            "clear_tilelang_kernel_caches",
+            "cuda_arch_key",
+            "lru_cache",
+            "torch"
+          ],
+          "namespace": "decode",
+          "sha256": "b582151a6d2bd345cc5f87415543370e027c8f943a4dfe57a18300f08560bd73"
+        },
+        "kernel_tilelang_state.py": {
+          "exports": [
+            "Any",
+            "BFLOAT16",
+            "CUDA_GET_DEVICE_CAPABILITY",
+            "EXACT_FUSED_STATE_PROJECTION_CAPABILITIES",
+            "FLOAT16",
+            "FLOAT32",
+            "IS_GRAD_ENABLED",
+            "TORCH_STACK",
+            "TORCH_WHERE",
+            "TORCH_ZEROS_LIKE",
+            "_compiled_backward_kernel",
+            "_compiled_fast_state_scan",
+            "_compiled_fused_state_update",
+            "_compiled_kernel",
+            "_compiled_low_precision_state_kernel",
+            "_compiled_post_state",
+            "_compiled_state_projection",
+            "_compiled_x_mix",
+            "_dtype_name",
+            "_require_exact_fused_state_projection",
+            "build_fast_state_scan_program",
+            "build_fused_state_update_program",
+            "build_low_precision_state_program",
+            "build_post_state_program",
+            "build_state_backward_program",
+            "build_state_program",
+            "build_state_projection_program",
+            "build_x_mix_program",
+            "clear_tilelang_state_kernel_caches",
+            "cuda_arch_key",
+            "exact_fused_state_projection_supported",
+            "lru_cache",
+            "tilelang_fast_state_scan",
+            "tilelang_fused_state_update",
+            "tilelang_low_precision_state_finalize",
+            "tilelang_post_state",
+            "tilelang_state_finalize",
+            "tilelang_state_finalize_backward",
+            "tilelang_state_projection",
+            "tilelang_state_scan",
+            "tilelang_state_update",
+            "tilelang_x_mix",
+            "torch",
+            "torch_state_scan",
+            "torch_state_update"
+          ],
+          "namespace": "state",
+          "sha256": "a3597160cb8d7317151200d05b241af7edcad2928c3731687764d914c28bf21a"
+        }
+      }
+    },
+    "transforms": {
+      "configuration_rwkv7.py": {
+        "output_sha256": "60a060199526f5a6c4d5a2436520eed97e9f21f2ea3cb6e6584d9b5fcb6f21f4",
+        "source_sha256": "60a060199526f5a6c4d5a2436520eed97e9f21f2ea3cb6e6584d9b5fcb6f21f4"
+      },
+      "custom_ops.py": {
+        "output_sha256": "8240d4eebc9decbd5ebf8fbcc3abc0ee6441b0eb4854b8d7bc18172cd2babf17",
+        "source_sha256": "fb86e5c9dfc9deb5a1be821ce02c3a390ce5c95ff55feded6add2422f0dc3623"
+      },
+      "kernel_dispatch.py": {
+        "output_sha256": "0290559437016ba4e2c2a26f7c7a0985b6fbab991535b7ae65ba18eb421fa9e1",
+        "source_sha256": "0290559437016ba4e2c2a26f7c7a0985b6fbab991535b7ae65ba18eb421fa9e1"
+      },
+      "modeling_rwkv7.py": {
+        "output_sha256": "ad7f8c981d21a3b0361cd26a2c68234a1bfbacc31fe52b49c3fe9d79e5b32844",
+        "source_sha256": "7e1e57a4516efe78a74f971ada0fe58cb71c87102479ed7516b6d3585964ca21"
+      },
+      "state.py": {
+        "output_sha256": "d13d81af6a301a70024494f03b95006bf5f2471ba0b57fd7f1f94bd54f70f2ef",
+        "source_sha256": "d13d81af6a301a70024494f03b95006bf5f2471ba0b57fd7f1f94bd54f70f2ef"
+      },
+      "tilelang_decode.py": {
+        "output_sha256": "313120b38178642c17ef16531c4510de87dc0f0d79d80f70b6ade8818b4dfbe7",
+        "source_sha256": "fe63e6332037cacbd0400d82f65fd6f1c282e31d92906c5f34d39eb9e83eb642"
+      }
+    }
+  },
+  "schema_version": 6,
+  "source": {
+    "filename": "rwkv7-g1i-1.5b-20260805-ctx16384.pth",
+    "kind": "huggingface",
+    "reference": "BlinkDL/rwkv7-g1/rwkv7-g1i-1.5b-20260805-ctx16384.pth",
+    "revision": "ede85bf8ab2e59aff7d7ca909fbbc73317866d89",
+    "sha256": "32ef7b5bf4dc8bde843cf26dfad809a1f527e2e76a9e790e7d406e71bcd785da",
+    "size_bytes": 3055444605
+  }
+}

tokenizer_config.json

@@ -0,0 +1,18 @@
+{
+  "added_tokens_decoder": {
+    "0": {
+      "content": "<|endoftext|>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false,
+      "special": true
+    }
+  },
+  "backend": "tokenizers",
+  "eos_token": "<|endoftext|>",
+  "pad_token": "<|endoftext|>",
+  "padding_side": "left",
+  "tokenizer_class": "PreTrainedTokenizerFast",
+  "unk_token": "<|endoftext|>"
+}