Duplicate from nvidia/LocateAnything-3B

Co-authored-by: Zhiding Yu <Zhiding@users.noreply.huggingface.co>

.gitattributes

@@ -0,0 +1,43 @@
+*.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
+teaser.jpg filter=lfs diff=lfs merge=lfs -text
+assets/decoding_demo.mp4 filter=lfs diff=lfs merge=lfs -text
+assets/demo.mp4 filter=lfs diff=lfs merge=lfs -text
+assets/coco_lvis.png filter=lfs diff=lfs merge=lfs -text
+assets/dense_object_detection.png filter=lfs diff=lfs merge=lfs -text
+assets/layout_ocr.png filter=lfs diff=lfs merge=lfs -text
+assets/referring.png filter=lfs diff=lfs merge=lfs -text
+assets/sspro.png filter=lfs diff=lfs merge=lfs -text

LICENSE

@@ -0,0 +1,74 @@
+                                   NVIDIA License
+
+1. Definitions
+
+“Licensor” means any person or entity that distributes its Work.
+
+“Work” means (a) the original work of authorship made available under this license, 
+which may include software, documentation, or other files, and (b) any additions to or 
+derivative works  thereof  that are made available under this license.
+
+The terms “reproduce,” “reproduction,” “derivative works,” and “distribution” have the 
+meaning as provided under U.S. copyright law; provided, however, that 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.
+
+Works are “made available” under this license by including in or with the Work either 
+(a) a copyright notice referencing the applicability of this license to the Work, or 
+(b) a copy of this license.
+
+2. License Grant
+
+2.1 Copyright Grant. Subject to the terms and conditions of this license, each Licensor 
+grants to you a perpetual, worldwide, non-exclusive, royalty-free, copyright license to 
+use, reproduce, prepare derivative works of, publicly display, publicly perform, 
+sublicense and distribute its Work and any resulting derivative works in any form.
+
+3. Limitations
+
+3.1 Redistribution. You may reproduce or distribute the Work only if (a) you do so under 
+this license, (b) you include a complete copy of this license with your distribution, 
+and (c) you retain without modification any copyright, patent, trademark, or attribution 
+notices that are present in the Work.
+
+3.2 Derivative Works. You may specify that additional or different terms apply to the 
+use, reproduction, and distribution of your derivative works of the Work (“Your Terms”) 
+only if (a) Your Terms provide that the use limitation in Section 3.3 applies to your 
+derivative works, and (b) you identify the specific derivative works that are subject 
+to Your Terms. Notwithstanding Your Terms, this license (including the redistribution 
+requirements in Section 3.1) will continue to apply to the Work itself.
+
+3.3 Use Limitation. The Work and any derivative works thereof only may be used or 
+intended for use non-commercially. Notwithstanding the foregoing, NVIDIA Corporation 
+and its affiliates may use the Work and any derivative works commercially. As used 
+herein, “non-commercially” means for research or evaluation purposes only.
+
+3.4 Patent Claims. If you bring or threaten to bring a patent claim against any 
+Licensor (including any claim, cross-claim or counterclaim in a lawsuit) to enforce any 
+patents that you allege are infringed by any Work, then your rights under this license 
+from such Licensor (including the grant in Section 2.1) will terminate immediately.
+
+3.5 Trademarks. This license does not grant any rights to use any Licensor’s or its 
+affiliates’ names, logos, or trademarks, except as necessary to reproduce the notices 
+described in this license.
+
+3.6 Termination. If you violate any term of this license, then your rights under this 
+license (including the grant in Section 2.1) will terminate immediately.
+
+4. Disclaimer of Warranty.
+
+THE WORK IS PROVIDED “AS IS” WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, EITHER 
+EXPRESS OR IMPLIED, INCLUDING WARRANTIES OR CONDITIONS OF MERCHANTABILITY, FITNESS FOR 
+A PARTICULAR PURPOSE, TITLE OR NON-INFRINGEMENT. YOU BEAR THE RISK OF UNDERTAKING ANY 
+ACTIVITIES UNDER THIS LICENSE. 
+
+5. Limitation of Liability.
+
+EXCEPT AS PROHIBITED BY APPLICABLE LAW, IN NO EVENT AND UNDER NO LEGAL THEORY, WHETHER 
+IN TORT (INCLUDING NEGLIGENCE), CONTRACT, OR OTHERWISE SHALL ANY LICENSOR BE LIABLE TO 
+YOU FOR DAMAGES, INCLUDING ANY DIRECT, INDIRECT, SPECIAL, INCIDENTAL, OR CONSEQUENTIAL 
+DAMAGES ARISING OUT OF OR RELATED TO THIS LICENSE, THE USE OR INABILITY TO USE THE WORK 
+(INCLUDING BUT NOT LIMITED TO LOSS OF GOODWILL, BUSINESS INTERRUPTION, LOST PROFITS OR 
+DATA, COMPUTER FAILURE OR MALFUNCTION, OR ANY OTHER DAMAGES OR LOSSES), EVEN IF THE 
+LICENSOR HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGES.
+

README.md

@@ -0,0 +1,556 @@
+---
+license: other
+license_name: nvidia-license
+license_link: https://huggingface.co/nvidia/LocateAnything-3B/blob/main/LICENSE
+language:
+- en
+tags:
+- nvidia
+- eagle
+- vision
+- object-detection
+- grounding
+- locateanything
+- arxiv:2605.27365
+demo: https://huggingface.co/spaces/nvidia/LocateAnything
+github: https://github.com/NVlabs/Eagle/tree/main/Embodied
+library_name: transformers
+pipeline_tag: image-text-to-text
+base_model:
+- Qwen/Qwen2.5-3B-Instruct
+---
+
+# LocateAnything: Fast and High-Quality Vision-Language Grounding with Parallel Box Decoding
+
+<p align="center">
+  <img src="assets/teaser.jpg" alt="LocateAnything teaser" width="100%">
+</p>
+
+## 🔗 Quick Links
+
+* 🚀 **Online Demo**: [LocateAnything (Hugging Face Spaces)](https://huggingface.co/spaces/nvidia/LocateAnything)
+* 💻 **GitHub Code**: [NVlabs/Eagle/Embodied](https://github.com/NVlabs/Eagle/tree/main/Embodied)
+* 📄 **Paper**: [arXiv:2605.27365](https://arxiv.org/abs/2605.27365)
+
+
+# Model Overview
+
+### Description:
+
+LocateAnything is a vision-language model for fast and high-quality visual grounding, enabling precise object localization, dense detection, and point-based localization across diverse domains in both Enterprise Intelligence and Physical AI. The model adopts a generalist design, supporting tasks such as referring expression grounding, multi-object detection, GUI element grounding, and text localization, with strong performance in complex and cluttered scenes.
+
+Its core innovation, Parallel Box Decoding (PBD), predicts complete bounding box coordinates in a single parallel step rather than autoregressive token-by-token decoding, improving efficiency while preserving geometric consistency. This enables up to 2.5× higher throughput compared to prior approaches.
+
+The model is trained on a large-scale multi-domain dataset (12M images, 138M+ queries, 785M bounding boxes) spanning natural scenes, robotics, driving, GUI interaction, and document understanding. It serves as a foundation for generalist multimodal perception and has been integrated into NVIDIA’s frontier production-grade vision-language models, such as Nemotron 3 Nano Omni, supporting grounding, GUI understanding, and multimodal agentic capabilities.
+
+LocateAnything is developed as part of the [Eagle VLM](https://github.com/NVlabs/EAGLE) model family. This released model is for research and development only. In addition, LocateAnything contributed to [Nemotron](https://www.nvidia.com/en-us/ai-data-science/foundation-models/nemotron/) and [Cosmos](https://www.nvidia.com/en-us/ai/cosmos/) as part of the Computer Use and Visual Grounding features. We give special thanks the Nemotron and Cosmos Teams for time and efforts in product integration.
+
+### Demo Videos
+
+<p align="left">
+  <video src="https://huggingface.co/nvidia/LocateAnything-3B/resolve/main/assets/demo.mp4" controls="controls" width="80%">
+    Your browser does not support the video tag.
+  </video>
+</p>
+
+<p align="left">
+  <video src="https://huggingface.co/nvidia/LocateAnything-3B/resolve/main/assets/decoding_demo.mp4" controls="controls" width="80%">
+    Your browser does not support the video tag.
+  </video>
+</p>
+
+### License/Terms of Use:
+
+This model is released under the [NVIDIA License](https://huggingface.co/nvidia/LocateAnything-3B/blob/main/LICENSE) for non-commercial use, which permits use, reproduction, and modification for **academic and non-profit research purposes only**. Commercial use is **not permitted**, except by NVIDIA and its affiliates. Redistribution must retain the license and all applicable copyright and attribution notices. The model is provided **“as is” without warranty of any kind**, and users assume all associated risks.
+
+This model is built using components from third-party models with their respective licenses:
+- Language model: [Qwen2.5-3B-Instruct](https://huggingface.co/Qwen/Qwen2.5-3B-Instruct) (Qwen Research License)
+- Vision encoder: [MoonViT-SO-400M](https://huggingface.co/moonshotai/MoonViT-SO-400M) (MIT License)
+
+Models are improved using Qwen.
+
+### Deployment Geography:
+
+Global
+
+### Use Case:
+
+LocateAnything-3B is intended for developers and researchers building vision-language models and applications that require fast and precise visual localization from natural language instructions.
+
+Supported use cases include:
+- Open-set, common, and long-tail object detection
+- Dense multi-object detection in cluttered scenes
+- Phrase and referring-expression grounding
+- Automated dataset labeling and annotation (e.g., detection, grounding, pointing)
+- GUI element grounding for interactive and agentic systems
+- Robotics and autonomous driving perception
+- Document understanding, layout grounding, and OCR localization
+- Industrial inspection, surveillance, and remote sensing applications
+- Point-based localization and fine-grained spatial reasoning
+
+### Release Date [Insert the expected release date below]:
+
+- Github [05/26/2026] via https://github.com/NVlabs/Eagle/tree/main/Embodied.
+- Hugging Face [05/26/2026] via https://huggingface.co/nvidia/LocateAnything-3B.
+- Demo [05/26/2026] via https://huggingface.co/spaces/nvidia/LocateAnything.
+- Webpage [05/26/2026] via https://research.nvidia.com/labs/lpr/locate-anything/.
+- Tech Report [05/26/2026] via https://research.nvidia.com/labs/lpr/locate-anything/LocateAnything.pdf
+
+## References(s):
+- Wang et al., [LocateAnything: Fast and High-Quality Vision-Language Grounding with Parallel Box Decoding](https://research.nvidia.com/labs/lpr/locate-anything/LocateAnything.pdf), NVIDIA Tech Report, 2026
+- Kimi Team, [Kimi-VL Technical Report](https://arxiv.org/abs/2504.07491), arXiv:2504.07491, 2025.
+- Qwen Team, [Qwen2.5: A Party of Foundation Models](https://qwen.ai/blog?id=qwen2.5), Qwen Blog, 2024.
+- Chen et al., [Pix2Seq: A Language Modeling Framework for Object Detection](https://arxiv.org/abs/2109.10852), ICLR, 2022.
+- Jiang et al., [Detect Anything via Next Point Prediction](https://arxiv.org/abs/2510.12798), arXiv:2510.12798, 2025.
+- Liu et al., [Grounding DINO: Marrying DINO with Grounded Pre-Training for Open-Set Object Detection](https://arxiv.org/abs/2303.05499), arXiv:2303.05499, 2023.
+- Lin et al., [Microsoft COCO: Common Objects in Context](https://arxiv.org/abs/1405.0312), ECCV, 2014.
+- Gupta et al., [LVIS: A Dataset for Large Vocabulary Instance Segmentation](https://arxiv.org/abs/1908.03195), CVPR, 2019.
+- Li et al., [ScreenSpot-Pro: GUI Grounding for Professional High-Resolution Computer Use](https://arxiv.org/abs/2504.07981), ACM MM, 2025.
+
+## Model Architecture:
+
+**Architecture Type:** Transformer-based vision-language model (VLM).  
+
+**Network Architecture:** Native-resolution VLM with the following components:
+- Vision encoder: MoonViT
+- Language model: Qwen2.5-3B-Instruct
+- Multimodal projector: MLP projector
+- Output formulation: Block-based structure for visual grounding
+
+**Number of model parameters:** 3B.
+
+LocateAnything extends a vision-language model with Parallel Box Decoding (PBD), a block-wise multi-token prediction framework for efficient visual grounding. Instead of autoregressive coordinate generation, the model predicts complete bounding boxes and points in parallel structured units, improving decoding efficiency while preserving geometric consistency. The architecture jointly optimizes next-token prediction and multi-token prediction to balance reasoning ability and parallel inference. Training follows a four-stage pipeline: initial multimodal knowledge adaptation using captioning, VQA, OCR, and related data, followed by grounding and dense-scene localization fine-tuning.
+
+## Input(s):
+
+**Input Type(s):** Image and Text.
+
+**Input Format(s):**
+- Image: RGB image input with original source resolution.
+- Text: Natural-language prompt or task template, such as object categories, referring expressions, GUI instructions, OCR/layout requests, or pointing queries.
+
+**Input Parameters:**
+- Image: Two-Dimensional (2D)
+- Text: One-Dimensional (1D)
+
+**Other Properties Related to Input:**
+- Production image resolution supports up to 2.5K.
+- Prompt length supports up to 24K tokens.
+- Training detection and grounding stages use a maximum sequence length of 25,600 tokens.
+- Inference supports up to 8,192 newly generated tokens.
+
+## Output(s):
+
+**Output Type(s):** Text.
+
+**Output Format(s):**
+- Text: Model-generated token sequence containing semantic labels and structured coordinate tokens, such as bounding boxes (`<box> x1, y1, x2, y2 </box>`) and points (`<box> x, y </box>`).
+
+**Output Parameters:**
+- Text: One-Dimensional (1D)
+- Bounding boxes/points: Two-Dimensional (2D) spatial coordinates
+
+**Other Properties Related to Output:**
+- Outputs are organized into fixed-length blocks (length 6), including Semantic, Box, Negative, and End blocks.
+- A Box block encodes quantized spatial coordinates with structural tokens; unused positions are padded with `<null>`.
+- Fast Mode predicts box-aligned blocks in parallel; Slow Mode uses autoregressive decoding; Hybrid Mode defaults to parallel decoding with fallback to autoregressive decoding for format irregularity or spatial ambiguity.
+
+Our AI models are designed and optimized to run on NVIDIA GPU-accelerated systems. By leveraging NVIDIA hardware (e.g., GPU cores) and software frameworks (e.g., CUDA libraries), the model achieves improved training and inference performance compared to CPU-only solutions.
+
+## Software Integration:
+**Runtime Engine(s):** 
+* Transformers. The inference setup uses standard VLM generation with BF16 precision and KV cache. TensorRT, TensorRT-LLM, and Triton are not yet supported.
+
+**Supported Hardware Microarchitecture Compatibility:**
+
+* NVIDIA Ampere (e.g., A100)
+* NVIDIA Blackwell
+* NVIDIA Hopper (e.g., H100)
+* NVIDIA Lovelace (e.g., L40, RTX 4090)
+
+Deployment on embedded platforms such as NVIDIA Thor is possible with additional model optimization, including quantization, compression, or distillation. Other architectures may be supported depending on available memory, precision support, and software configuration.
+
+**Supported Operating System(s):**
+* Linux
+
+The integration of foundation and fine-tuned models into AI systems requires additional testing using use-case-specific data to ensure safe and effective deployment. Following the V-model methodology, iterative testing and validation at both unit and system levels are essential to mitigate risks, meet technical and functional requirements, and ensure compliance with safety and ethical standards before deployment.
+
+## Model Version(s):
+LocateAnything-3B: 3B-parameter research model variant evaluated in Hybrid Mode by default. Fast, Hybrid, and Slow inference modes are supported by the same model formulation.
+
+LocateAnything-3B can be integrated into systems that require spatial grounding from natural language, such as GUI agents, robotics/embodied agents, document-understanding pipelines, OCR/text localization, and open-world detection workflows.
+
+## Training, Testing, and Evaluation Datasets:
+
+### Data Modality:
+Image and Text. <br>
+* Image <br>
+* Text <br>
+
+### Training Data Size: 
+**Image Training Data Size:** <br>
+* 1 Million to 1 Billion Images - 12M unique images. <br>
+
+**Text Training Data Size:** <br>
+* 1 Billion to 10 Trillion Tokens - Derived from approximately 140M natural-language queries. <br>
+
+**Data Collection Method by dataset:** <br>
+- Hybrid: Human, Automated <br>
+  Data is collected from human-curated and open-source datasets, as well as automated ingestion of publicly available data sources.
+
+**Labeling Method by dataset:** <br>
+- Hybrid: Human, Synthetic, Automated <br>
+  Labeling includes original human or open-source annotations, along with model-assisted and synthetic annotation generation using Qwen3-VL, Molmo, SAM 3, and Rex-Omni, with automated post-verification.
+
+**Properties:** The training data consists of supervised fine-tuning (SFT) datasets with multimodal inputs, primarily image-text pairs and structured annotations such as bounding boxes, points, and negative samples.
+
+The data spans multiple domains, including grounding, open-world grounding, general and dense object detection, scene text detection, GUI understanding and grounding, document layout understanding, and OCR.
+
+Modalities include visual inputs (images) and natural-language queries or instructions. The dataset is derived from a mixture of publicly available academic datasets, along with model-assisted and synthetic annotations. It may include publicly available and potentially copyrighted content; users are responsible for ensuring compliance with applicable usage rights.
+
+The linguistic content primarily consists of short, task-oriented natural-language expressions, such as object categories, referring expressions, GUI instructions, OCR queries, and grounding prompts, typically in English.
+
+## Evaluation Dataset:
+
+**Data Collection Method by dataset:**
+- Hybrid: Human, Automated
+
+**Labeling Method by dataset:**
+- Hybrid: Human, Synthetic, Automated
+
+**Properties:** The evaluation datasets consist of publicly available benchmarks spanning visual grounding, object detection, document understanding, scene text detection, and GUI-related tasks. Modalities include image inputs paired with natural-language queries and structured annotations such as bounding boxes and points.
+
+The evaluation suite covers both box-level and point-level grounding tasks, with approximately 48K images for box evaluation and 35K images for point evaluation across multiple datasets. These datasets span diverse domains including natural scenes, documents, aerial imagery, and human-centric interactions, enabling comprehensive assessment of localization accuracy and robustness.
+
+Evaluation queries are typically short, task-oriented natural-language expressions such as referring phrases, object categories, and grounding prompts.
+
+Performance is measured using box-based F1 at IoU thresholds of 0.5 and 0.95, as well as mean IoU for detection, layout, and OCR tasks. Point-based localization is evaluated based on whether predicted points fall within ground-truth segmentation masks or bounding boxes. Inference efficiency is reported in boxes per second (BPS) on a single NVIDIA H100 GPU with batch size 1.
+
+## Quantitative Evaluation Benchmarks
+
+### General Object Detection
+<p align="left">
+  <img src="assets/coco_lvis.png" width="700">
+</p>
+
+### Dense Object Detection
+<p align="left">
+  <img src="assets/dense_object_detection.png" width="700">
+</p>
+
+### GUI Understanding
+<p align="left">
+  <img src="assets/sspro.png" width="700">
+</p>
+
+### Layout Grounding and OCR
+<p align="left">
+  <img src="assets/layout_ocr.png" width="700">
+</p>
+
+### Referring Expression Grounding
+<p align="left">
+  <img src="assets/referring.png" width="700">
+</p>
+
+### Pointing
+<p align="left">
+  <img src="assets/pointing.png" width="700">
+</p>
+
+## Inference:
+
+Test Hardware: H100 & A100
+
+We suggest using `max_new_tokens=8192` and `generation_mode="hybrid"` to avoid truncated response and balance speed with accuracy.
+
+### Batch Hybrid Inference
+
+This release includes `batch_infer.py`, `batch_utils`, and `kernel_utils` for
+high-throughput detection and grounding. The `la_flash` backend is a pure
+FlashAttention-varlen sparse range executor: it keeps LocateAnything's hybrid
+MTP decoding path, avoids dense `[B,H,Q,K]` SDPA masks, and does not require a
+custom CUDA extension build.
+
+Use it with:
+
+```bash
+python batch_infer.py \
+  --model . \
+  --attn la_flash \
+  --scheduler pipeline \
+  --batch-size 4 \
+  --image /path/to/image.jpg \
+  --query "person</c>car"
+```
+
+A100 4K probe, real 3840x2160 street image, `query=vehicle`,
+`batch_size=4`, raw PIL input, `in_token_limit=25600`, hybrid MTP inference:
+
+| Backend | Attention Path | Time | Peak Reserved Memory |
+| --- | --- | ---: | ---: |
+| `sdpa` | Dense SDPA masks | 8.2600 s | 35.12 GB |
+| `la_flash` | FlashAttention sparse range plan | 8.0314 s | 11.71 GB |
+
+See `batch_utils/README.md` and `kernel_utils/README.md` for runtime knobs and
+implementation details.
+
+### Installation
+
+```bash
+pip install opencv-python-headless==4.11.0.86 transformers==4.57.1 numpy==1.25.0 Pillow==11.1.0 peft torchvision decord==0.6.0 lmdb==1.7.5
+```
+
+> PyTorch (`torch`) must be installed separately according to your CUDA version. See [pytorch.org/get-started](https://pytorch.org/get-started/locally/).
+
+Optional — [MagiAttention](https://sandai-org.github.io/MagiAttention/docs/main/user_guide/install.html) (Hopper / Blackwell GPUs only, recommended for faster MTP inference):
+
+```bash
+git clone https://github.com/SandAI-org/MagiAttention.git
+cd MagiAttention
+git checkout v1.0.5
+git submodule update --init --recursive
+pip install -r requirements.txt
+pip install --no-build-isolation .
+```
+
+If MagiAttention is installed, the model will automatically use it for efficient MTP block-diffusion attention. If not installed, it will fall back to PyTorch SDPA — fully functional but slower for MTP decoding.
+
+### Worker (recommended)
+
+Below is a self-contained worker that loads the model once and serves perception queries via a unified `predict()` plus task-specific convenience methods. You can drop this class into any FastAPI / gRPC / Triton serving framework.
+
+```python
+import re
+import torch
+from PIL import Image
+from transformers import AutoModel, AutoTokenizer, AutoProcessor
+
+
+class LocateAnythingWorker:
+    """Stateful worker that loads the model once and serves perception queries."""
+
+    def __init__(self, model_path: str, device: str = "cuda", dtype=torch.bfloat16):
+        self.device = device
+        self.dtype = dtype
+
+        self.tokenizer = AutoTokenizer.from_pretrained(model_path, trust_remote_code=True)
+        self.processor = AutoProcessor.from_pretrained(model_path, trust_remote_code=True)
+        self.model = AutoModel.from_pretrained(
+            model_path,
+            torch_dtype=dtype,
+            trust_remote_code=True,
+        ).to(device).eval()
+
+    @torch.no_grad()
+    def predict(
+        self,
+        image: Image.Image,
+        question: str,
+        generation_mode: str = "hybrid",   # "fast" (MTP) | "slow" (NTP/AR) | "hybrid"
+        max_new_tokens: int = 2048,
+        temperature: float = 0.7,
+        verbose: bool = True,
+    ) -> dict:
+        messages = [
+            {"role": "user", "content": [
+                {"type": "image", "image": image},
+                {"type": "text", "text": question},
+            ]}
+        ]
+
+        text = self.processor.py_apply_chat_template(
+            messages, tokenize=False, add_generation_prompt=True
+        )
+        images, videos = self.processor.process_vision_info(messages)
+        inputs = self.processor(
+            text=[text], images=images, videos=videos, return_tensors="pt"
+        ).to(self.device)
+
+        pixel_values = inputs["pixel_values"].to(self.dtype)
+        input_ids = inputs["input_ids"]
+        image_grid_hws = inputs.get("image_grid_hws", None)
+
+        response = self.model.generate(
+            pixel_values=pixel_values,
+            input_ids=input_ids,
+            attention_mask=inputs["attention_mask"],
+            image_grid_hws=image_grid_hws,
+            tokenizer=self.tokenizer,
+            max_new_tokens=max_new_tokens,
+            use_cache=True,
+            generation_mode=generation_mode,
+            temperature=temperature,
+            do_sample=True,
+            top_p=0.9,
+            repetition_penalty=1.1,
+            verbose=verbose,
+        )
+
+        result = {"answer": response[0] if isinstance(response, tuple) else response}
+        if isinstance(response, tuple) and len(response) >= 3:
+            result["history"] = response[1]
+            result["stats"] = response[2]
+        return result
+
+    # ---- Convenience methods for each task ----
+
+    def detect(self, image: Image.Image, categories: list[str], **kwargs) -> dict:
+        """Object detection / document layout analysis."""
+        cats = "</c>".join(categories)
+        prompt = f"Locate all the instances that matches the following description: {cats}."
+        return self.predict(image, prompt, **kwargs)
+
+    def ground_single(self, image: Image.Image, phrase: str, **kwargs) -> dict:
+        """Phrase grounding — single instance."""
+        prompt = f"Locate a single instance that matches the following description: {phrase}."
+        return self.predict(image, prompt, **kwargs)
+
+    def ground_multi(self, image: Image.Image, phrase: str, **kwargs) -> dict:
+        """Phrase grounding — multiple instances."""
+        prompt = f"Locate all the instances that match the following description: {phrase}."
+        return self.predict(image, prompt, **kwargs)
+
+    def ground_text(self, image: Image.Image, phrase: str, **kwargs) -> dict:
+        """Text grounding."""
+        prompt = f"Please locate the text referred as {phrase}."
+        return self.predict(image, prompt, **kwargs)
+
+    def detect_text(self, image: Image.Image, **kwargs) -> dict:
+        """Scene text detection."""
+        prompt = "Detect all the text in box format."
+        return self.predict(image, prompt, **kwargs)
+
+    def ground_gui(self, image: Image.Image, phrase: str, output_type: str = "box", **kwargs) -> dict:
+        """GUI grounding (box or point)."""
+        if output_type == "point":
+            prompt = f"Point to: {phrase}."
+        else:
+            prompt = f"Locate the region that matches the following description: {phrase}."
+        return self.predict(image, prompt, **kwargs)
+
+    def point(self, image: Image.Image, phrase: str, **kwargs) -> dict:
+        """Pointing."""
+        prompt = f"Point to: {phrase}."
+        return self.predict(image, prompt, **kwargs)
+
+    # ---- Utility: parse model output ----
+
+    @staticmethod
+    def parse_boxes(answer: str, image_width: int, image_height: int) -> list[dict]:
+        """Parse model output into pixel-coordinate bounding boxes.
+
+        Coordinates in model output are normalized integers in [0, 1000].
+        """
+        boxes = []
+        for m in re.finditer(r"<box><(\d+)><(\d+)><(\d+)><(\d+)></box>", answer):
+            x1, y1, x2, y2 = [int(g) for g in m.groups()]
+            boxes.append({
+                "x1": x1 / 1000 * image_width,
+                "y1": y1 / 1000 * image_height,
+                "x2": x2 / 1000 * image_width,
+                "y2": y2 / 1000 * image_height,
+            })
+        return boxes
+
+    @staticmethod
+    def parse_points(answer: str, image_width: int, image_height: int) -> list[dict]:
+        """Parse model output into pixel-coordinate points."""
+        points = []
+        for m in re.finditer(r"<box><(\d+)><(\d+)></box>", answer):
+            x, y = int(m.group(1)), int(m.group(2))
+            points.append({
+                "x": x / 1000 * image_width,
+                "y": y / 1000 * image_height,
+            })
+        return points
+```
+
+### Usage Example
+
+```python
+worker = LocateAnythingWorker("nvidia/LocateAnything-3B")
+img = Image.open("example.jpg").convert("RGB")
+
+# Object Detection
+result = worker.detect(img, ["person", "car", "bicycle"])
+print("Detection:", result["answer"])
+
+# Phrase Grounding (multiple)
+result = worker.ground_multi(img, "people wearing red shirts")
+print("Grounding:", result["answer"])
+
+# Scene Text Detection
+result = worker.detect_text(img)
+print("Text Detection:", result["answer"])
+
+# Pointing
+result = worker.point(img, "the traffic light")
+print("Pointing:", result["answer"])
+
+# GUI Grounding (point)
+result = worker.ground_gui(img, "the search button", output_type="point")
+print("GUI Point:", result["answer"])
+
+# Parse structured output into pixel coordinates
+w, h = img.size
+boxes = LocateAnythingWorker.parse_boxes(result["answer"], w, h)
+points = LocateAnythingWorker.parse_points(result["answer"], w, h)
+```
+
+### Supported Tasks & Prompt Templates
+
+| Task | Worker Method | Output | Prompt Template |
+| --- | --- | --- | --- |
+| Object Detection | `worker.detect(img, [...])` | Box | `Locate all the instances that matches the following description: [CATEGORIES].` |
+| Phrase Grounding (single) | `worker.ground_single(img, phrase)` | Single Box | `Locate a single instance that matches the following description: [PHRASE].` |
+| Phrase Grounding (multi) | `worker.ground_multi(img, phrase)` | Multiple Boxes | `Locate all the instances that match the following description: [PHRASE].` |
+| Text Grounding | `worker.ground_text(img, phrase)` | Box | `Please locate the text referred as [PHRASE].` |
+| Scene Text Detection | `worker.detect_text(img)` | Box | `Detect all the text in box format.` |
+| Document Layout Analysis | `worker.detect(img, [...])` | Box | `Locate all the instances that matches the following description: [CATEGORIES].` |
+| GUI Grounding (box) | `worker.ground_gui(img, phrase, "box")` | Box | `Locate the region that matches the following description: [PHRASE].` |
+| GUI Grounding (point) / Pointing | `worker.ground_gui(img, phrase, "point")` / `worker.point(img, phrase)` | Point | `Point to: [PHRASE].` |
+
+`[PHRASE]` is a free-form natural-language description; `[CATEGORIES]` is a comma-separated list (multiple categories may also be joined with `</c>`).
+
+### Generation Modes
+
+| Mode | Description | Speed | Accuracy |
+| --- | --- | --- | --- |
+| `fast` | MTP only, never falls back to AR | Fastest | Good for simple scenes |
+| `slow` | Pure auto-regressive decoding | Slowest | Most robust |
+| `hybrid` (default) | MTP first, falls back to AR on uncertain boxes, switches back after box boundary | Balanced | Best overall |
+
+## Batch Utils and Kernel Utils
+
+This repository also includes optional utilities for high-throughput detection
+runs:
+
+- `batch_infer.py`: JSONL/image-query batch inference CLI.
+- `batch_utils/`: batched hybrid generation runtime. See
+  `batch_utils/README.md`.
+- `kernel_utils/`: LA Flash sparse range utilities. See
+  `kernel_utils/README.md`.
+
+Run a small batch inference job:
+
+```bash
+python batch_infer.py \
+  --model . \
+  --attn la_flash \
+  --scheduler pipeline \
+  --batch-size 4 \
+  --image assets/pointing.png \
+  --query "the object being pointed at"
+```
+
+The batched sparse-plan decode runtime is intended for inference/evaluation and
+does not support the training `labels` path. Training remains on the
+MagiAttention backend.
+
+## Ethical Considerations:
+NVIDIA believes Trustworthy AI is a shared responsibility and we have established policies and practices to enable development for a wide array of AI applications. When downloaded or used in accordance with our terms of service, developers should work with their internal model team to ensure this model meets requirements for the relevant industry and use case and addresses unforeseen product misuse.
+
+Please make sure you have proper rights and permissions for all input image and video content; if image or video includes people, personal health information, or intellectual property, the image or video generated will not blur or maintain proportions of image subjects included.
+
+Please report model quality, risk, security vulnerabilities or NVIDIA AI Concerns [here](https://app.intigriti.com/programs/nvidia/nvidiavdp/detail).

added_tokens.json

@@ -0,0 +1,1040 @@
+{
+  "</box>": 151669,
+  "</c>": 152680,
+  "</img>": 151667,
+  "</interval>": 151675,
+  "</quad>": 151671,
+  "</ref>": 151673,
+  "</tool_call>": 151658,
+  "<0>": 151677,
+  "<1000>": 152677,
+  "<100>": 151777,
+  "<101>": 151778,
+  "<102>": 151779,
+  "<103>": 151780,
+  "<104>": 151781,
+  "<105>": 151782,
+  "<106>": 151783,
+  "<107>": 151784,
+  "<108>": 151785,
+  "<109>": 151786,
+  "<10>": 151687,
+  "<110>": 151787,
+  "<111>": 151788,
+  "<112>": 151789,
+  "<113>": 151790,
+  "<114>": 151791,
+  "<115>": 151792,
+  "<116>": 151793,
+  "<117>": 151794,
+  "<118>": 151795,
+  "<119>": 151796,
+  "<11>": 151688,
+  "<120>": 151797,
+  "<121>": 151798,
+  "<122>": 151799,
+  "<123>": 151800,
+  "<124>": 151801,
+  "<125>": 151802,
+  "<126>": 151803,
+  "<127>": 151804,
+  "<128>": 151805,
+  "<129>": 151806,
+  "<12>": 151689,
+  "<130>": 151807,
+  "<131>": 151808,
+  "<132>": 151809,
+  "<133>": 151810,
+  "<134>": 151811,
+  "<135>": 151812,
+  "<136>": 151813,
+  "<137>": 151814,
+  "<138>": 151815,
+  "<139>": 151816,
+  "<13>": 151690,
+  "<140>": 151817,
+  "<141>": 151818,
+  "<142>": 151819,
+  "<143>": 151820,
+  "<144>": 151821,
+  "<145>": 151822,
+  "<146>": 151823,
+  "<147>": 151824,
+  "<148>": 151825,
+  "<149>": 151826,
+  "<14>": 151691,
+  "<150>": 151827,
+  "<151>": 151828,
+  "<152>": 151829,
+  "<153>": 151830,
+  "<154>": 151831,
+  "<155>": 151832,
+  "<156>": 151833,
+  "<157>": 151834,
+  "<158>": 151835,
+  "<159>": 151836,
+  "<15>": 151692,
+  "<160>": 151837,
+  "<161>": 151838,
+  "<162>": 151839,
+  "<163>": 151840,
+  "<164>": 151841,
+  "<165>": 151842,
+  "<166>": 151843,
+  "<167>": 151844,
+  "<168>": 151845,
+  "<169>": 151846,
+  "<16>": 151693,
+  "<170>": 151847,
+  "<171>": 151848,
+  "<172>": 151849,
+  "<173>": 151850,
+  "<174>": 151851,
+  "<175>": 151852,
+  "<176>": 151853,
+  "<177>": 151854,
+  "<178>": 151855,
+  "<179>": 151856,
+  "<17>": 151694,
+  "<180>": 151857,
+  "<181>": 151858,
+  "<182>": 151859,
+  "<183>": 151860,
+  "<184>": 151861,
+  "<185>": 151862,
+  "<186>": 151863,
+  "<187>": 151864,
+  "<188>": 151865,
+  "<189>": 151866,
+  "<18>": 151695,
+  "<190>": 151867,
+  "<191>": 151868,
+  "<192>": 151869,
+  "<193>": 151870,
+  "<194>": 151871,
+  "<195>": 151872,
+  "<196>": 151873,
+  "<197>": 151874,
+  "<198>": 151875,
+  "<199>": 151876,
+  "<19>": 151696,
+  "<1>": 151678,
+  "<200>": 151877,
+  "<201>": 151878,
+  "<202>": 151879,
+  "<203>": 151880,
+  "<204>": 151881,
+  "<205>": 151882,
+  "<206>": 151883,
+  "<207>": 151884,
+  "<208>": 151885,
+  "<209>": 151886,
+  "<20>": 151697,
+  "<210>": 151887,
+  "<211>": 151888,
+  "<212>": 151889,
+  "<213>": 151890,
+  "<214>": 151891,
+  "<215>": 151892,
+  "<216>": 151893,
+  "<217>": 151894,
+  "<218>": 151895,
+  "<219>": 151896,
+  "<21>": 151698,
+  "<220>": 151897,
+  "<221>": 151898,
+  "<222>": 151899,
+  "<223>": 151900,
+  "<224>": 151901,
+  "<225>": 151902,
+  "<226>": 151903,
+  "<227>": 151904,
+  "<228>": 151905,
+  "<229>": 151906,
+  "<22>": 151699,
+  "<230>": 151907,
+  "<231>": 151908,
+  "<232>": 151909,
+  "<233>": 151910,
+  "<234>": 151911,
+  "<235>": 151912,
+  "<236>": 151913,
+  "<237>": 151914,
+  "<238>": 151915,
+  "<239>": 151916,
+  "<23>": 151700,
+  "<240>": 151917,
+  "<241>": 151918,
+  "<242>": 151919,
+  "<243>": 151920,
+  "<244>": 151921,
+  "<245>": 151922,
+  "<246>": 151923,
+  "<247>": 151924,
+  "<248>": 151925,
+  "<249>": 151926,
+  "<24>": 151701,
+  "<250>": 151927,
+  "<251>": 151928,
+  "<252>": 151929,
+  "<253>": 151930,
+  "<254>": 151931,
+  "<255>": 151932,
+  "<256>": 151933,
+  "<257>": 151934,
+  "<258>": 151935,
+  "<259>": 151936,
+  "<25>": 151702,
+  "<260>": 151937,
+  "<261>": 151938,
+  "<262>": 151939,
+  "<263>": 151940,
+  "<264>": 151941,
+  "<265>": 151942,
+  "<266>": 151943,
+  "<267>": 151944,
+  "<268>": 151945,
+  "<269>": 151946,
+  "<26>": 151703,
+  "<270>": 151947,
+  "<271>": 151948,
+  "<272>": 151949,
+  "<273>": 151950,
+  "<274>": 151951,
+  "<275>": 151952,
+  "<276>": 151953,
+  "<277>": 151954,
+  "<278>": 151955,
+  "<279>": 151956,
+  "<27>": 151704,
+  "<280>": 151957,
+  "<281>": 151958,
+  "<282>": 151959,
+  "<283>": 151960,
+  "<284>": 151961,
+  "<285>": 151962,
+  "<286>": 151963,
+  "<287>": 151964,
+  "<288>": 151965,
+  "<289>": 151966,
+  "<28>": 151705,
+  "<290>": 151967,
+  "<291>": 151968,
+  "<292>": 151969,
+  "<293>": 151970,
+  "<294>": 151971,
+  "<295>": 151972,
+  "<296>": 151973,
+  "<297>": 151974,
+  "<298>": 151975,
+  "<299>": 151976,
+  "<29>": 151706,
+  "<2>": 151679,
+  "<300>": 151977,
+  "<301>": 151978,
+  "<302>": 151979,
+  "<303>": 151980,
+  "<304>": 151981,
+  "<305>": 151982,
+  "<306>": 151983,
+  "<307>": 151984,
+  "<308>": 151985,
+  "<309>": 151986,
+  "<30>": 151707,
+  "<310>": 151987,
+  "<311>": 151988,
+  "<312>": 151989,
+  "<313>": 151990,
+  "<314>": 151991,
+  "<315>": 151992,
+  "<316>": 151993,
+  "<317>": 151994,
+  "<318>": 151995,
+  "<319>": 151996,
+  "<31>": 151708,
+  "<320>": 151997,
+  "<321>": 151998,
+  "<322>": 151999,
+  "<323>": 152000,
+  "<324>": 152001,
+  "<325>": 152002,
+  "<326>": 152003,
+  "<327>": 152004,
+  "<328>": 152005,
+  "<329>": 152006,
+  "<32>": 151709,
+  "<330>": 152007,
+  "<331>": 152008,
+  "<332>": 152009,
+  "<333>": 152010,
+  "<334>": 152011,
+  "<335>": 152012,
+  "<336>": 152013,
+  "<337>": 152014,
+  "<338>": 152015,
+  "<339>": 152016,
+  "<33>": 151710,
+  "<340>": 152017,
+  "<341>": 152018,
+  "<342>": 152019,
+  "<343>": 152020,
+  "<344>": 152021,
+  "<345>": 152022,
+  "<346>": 152023,
+  "<347>": 152024,
+  "<348>": 152025,
+  "<349>": 152026,
+  "<34>": 151711,
+  "<350>": 152027,
+  "<351>": 152028,
+  "<352>": 152029,
+  "<353>": 152030,
+  "<354>": 152031,
+  "<355>": 152032,
+  "<356>": 152033,
+  "<357>": 152034,
+  "<358>": 152035,
+  "<359>": 152036,
+  "<35>": 151712,
+  "<360>": 152037,
+  "<361>": 152038,
+  "<362>": 152039,
+  "<363>": 152040,
+  "<364>": 152041,
+  "<365>": 152042,
+  "<366>": 152043,
+  "<367>": 152044,
+  "<368>": 152045,
+  "<369>": 152046,
+  "<36>": 151713,
+  "<370>": 152047,
+  "<371>": 152048,
+  "<372>": 152049,
+  "<373>": 152050,
+  "<374>": 152051,
+  "<375>": 152052,
+  "<376>": 152053,
+  "<377>": 152054,
+  "<378>": 152055,
+  "<379>": 152056,
+  "<37>": 151714,
+  "<380>": 152057,
+  "<381>": 152058,
+  "<382>": 152059,
+  "<383>": 152060,
+  "<384>": 152061,
+  "<385>": 152062,
+  "<386>": 152063,
+  "<387>": 152064,
+  "<388>": 152065,
+  "<389>": 152066,
+  "<38>": 151715,
+  "<390>": 152067,
+  "<391>": 152068,
+  "<392>": 152069,
+  "<393>": 152070,
+  "<394>": 152071,
+  "<395>": 152072,
+  "<396>": 152073,
+  "<397>": 152074,
+  "<398>": 152075,
+  "<399>": 152076,
+  "<39>": 151716,
+  "<3>": 151680,
+  "<400>": 152077,
+  "<401>": 152078,
+  "<402>": 152079,
+  "<403>": 152080,
+  "<404>": 152081,
+  "<405>": 152082,
+  "<406>": 152083,
+  "<407>": 152084,
+  "<408>": 152085,
+  "<409>": 152086,
+  "<40>": 151717,
+  "<410>": 152087,
+  "<411>": 152088,
+  "<412>": 152089,
+  "<413>": 152090,
+  "<414>": 152091,
+  "<415>": 152092,
+  "<416>": 152093,
+  "<417>": 152094,
+  "<418>": 152095,
+  "<419>": 152096,
+  "<41>": 151718,
+  "<420>": 152097,
+  "<421>": 152098,
+  "<422>": 152099,
+  "<423>": 152100,
+  "<424>": 152101,
+  "<425>": 152102,
+  "<426>": 152103,
+  "<427>": 152104,
+  "<428>": 152105,
+  "<429>": 152106,
+  "<42>": 151719,
+  "<430>": 152107,
+  "<431>": 152108,
+  "<432>": 152109,
+  "<433>": 152110,
+  "<434>": 152111,
+  "<435>": 152112,
+  "<436>": 152113,
+  "<437>": 152114,
+  "<438>": 152115,
+  "<439>": 152116,
+  "<43>": 151720,
+  "<440>": 152117,
+  "<441>": 152118,
+  "<442>": 152119,
+  "<443>": 152120,
+  "<444>": 152121,
+  "<445>": 152122,
+  "<446>": 152123,
+  "<447>": 152124,
+  "<448>": 152125,
+  "<449>": 152126,
+  "<44>": 151721,
+  "<450>": 152127,
+  "<451>": 152128,
+  "<452>": 152129,
+  "<453>": 152130,
+  "<454>": 152131,
+  "<455>": 152132,
+  "<456>": 152133,
+  "<457>": 152134,
+  "<458>": 152135,
+  "<459>": 152136,
+  "<45>": 151722,
+  "<460>": 152137,
+  "<461>": 152138,
+  "<462>": 152139,
+  "<463>": 152140,
+  "<464>": 152141,
+  "<465>": 152142,
+  "<466>": 152143,
+  "<467>": 152144,
+  "<468>": 152145,
+  "<469>": 152146,
+  "<46>": 151723,
+  "<470>": 152147,
+  "<471>": 152148,
+  "<472>": 152149,
+  "<473>": 152150,
+  "<474>": 152151,
+  "<475>": 152152,
+  "<476>": 152153,
+  "<477>": 152154,
+  "<478>": 152155,
+  "<479>": 152156,
+  "<47>": 151724,
+  "<480>": 152157,
+  "<481>": 152158,
+  "<482>": 152159,
+  "<483>": 152160,
+  "<484>": 152161,
+  "<485>": 152162,
+  "<486>": 152163,
+  "<487>": 152164,
+  "<488>": 152165,
+  "<489>": 152166,
+  "<48>": 151725,
+  "<490>": 152167,
+  "<491>": 152168,
+  "<492>": 152169,
+  "<493>": 152170,
+  "<494>": 152171,
+  "<495>": 152172,
+  "<496>": 152173,
+  "<497>": 152174,
+  "<498>": 152175,
+  "<499>": 152176,
+  "<49>": 151726,
+  "<4>": 151681,
+  "<500>": 152177,
+  "<501>": 152178,
+  "<502>": 152179,
+  "<503>": 152180,
+  "<504>": 152181,
+  "<505>": 152182,
+  "<506>": 152183,
+  "<507>": 152184,
+  "<508>": 152185,
+  "<509>": 152186,
+  "<50>": 151727,
+  "<510>": 152187,
+  "<511>": 152188,
+  "<512>": 152189,
+  "<513>": 152190,
+  "<514>": 152191,
+  "<515>": 152192,
+  "<516>": 152193,
+  "<517>": 152194,
+  "<518>": 152195,
+  "<519>": 152196,
+  "<51>": 151728,
+  "<520>": 152197,
+  "<521>": 152198,
+  "<522>": 152199,
+  "<523>": 152200,
+  "<524>": 152201,
+  "<525>": 152202,
+  "<526>": 152203,
+  "<527>": 152204,
+  "<528>": 152205,
+  "<529>": 152206,
+  "<52>": 151729,
+  "<530>": 152207,
+  "<531>": 152208,
+  "<532>": 152209,
+  "<533>": 152210,
+  "<534>": 152211,
+  "<535>": 152212,
+  "<536>": 152213,
+  "<537>": 152214,
+  "<538>": 152215,
+  "<539>": 152216,
+  "<53>": 151730,
+  "<540>": 152217,
+  "<541>": 152218,
+  "<542>": 152219,
+  "<543>": 152220,
+  "<544>": 152221,
+  "<545>": 152222,
+  "<546>": 152223,
+  "<547>": 152224,
+  "<548>": 152225,
+  "<549>": 152226,
+  "<54>": 151731,
+  "<550>": 152227,
+  "<551>": 152228,
+  "<552>": 152229,
+  "<553>": 152230,
+  "<554>": 152231,
+  "<555>": 152232,
+  "<556>": 152233,
+  "<557>": 152234,
+  "<558>": 152235,
+  "<559>": 152236,
+  "<55>": 151732,
+  "<560>": 152237,
+  "<561>": 152238,
+  "<562>": 152239,
+  "<563>": 152240,
+  "<564>": 152241,
+  "<565>": 152242,
+  "<566>": 152243,
+  "<567>": 152244,
+  "<568>": 152245,
+  "<569>": 152246,
+  "<56>": 151733,
+  "<570>": 152247,
+  "<571>": 152248,
+  "<572>": 152249,
+  "<573>": 152250,
+  "<574>": 152251,
+  "<575>": 152252,
+  "<576>": 152253,
+  "<577>": 152254,
+  "<578>": 152255,
+  "<579>": 152256,
+  "<57>": 151734,
+  "<580>": 152257,
+  "<581>": 152258,
+  "<582>": 152259,
+  "<583>": 152260,
+  "<584>": 152261,
+  "<585>": 152262,
+  "<586>": 152263,
+  "<587>": 152264,
+  "<588>": 152265,
+  "<589>": 152266,
+  "<58>": 151735,
+  "<590>": 152267,
+  "<591>": 152268,
+  "<592>": 152269,
+  "<593>": 152270,
+  "<594>": 152271,
+  "<595>": 152272,
+  "<596>": 152273,
+  "<597>": 152274,
+  "<598>": 152275,
+  "<599>": 152276,
+  "<59>": 151736,
+  "<5>": 151682,
+  "<600>": 152277,
+  "<601>": 152278,
+  "<602>": 152279,
+  "<603>": 152280,
+  "<604>": 152281,
+  "<605>": 152282,
+  "<606>": 152283,
+  "<607>": 152284,
+  "<608>": 152285,
+  "<609>": 152286,
+  "<60>": 151737,
+  "<610>": 152287,
+  "<611>": 152288,
+  "<612>": 152289,
+  "<613>": 152290,
+  "<614>": 152291,
+  "<615>": 152292,
+  "<616>": 152293,
+  "<617>": 152294,
+  "<618>": 152295,
+  "<619>": 152296,
+  "<61>": 151738,
+  "<620>": 152297,
+  "<621>": 152298,
+  "<622>": 152299,
+  "<623>": 152300,
+  "<624>": 152301,
+  "<625>": 152302,
+  "<626>": 152303,
+  "<627>": 152304,
+  "<628>": 152305,
+  "<629>": 152306,
+  "<62>": 151739,
+  "<630>": 152307,
+  "<631>": 152308,
+  "<632>": 152309,
+  "<633>": 152310,
+  "<634>": 152311,
+  "<635>": 152312,
+  "<636>": 152313,
+  "<637>": 152314,
+  "<638>": 152315,
+  "<639>": 152316,
+  "<63>": 151740,
+  "<640>": 152317,
+  "<641>": 152318,
+  "<642>": 152319,
+  "<643>": 152320,
+  "<644>": 152321,
+  "<645>": 152322,
+  "<646>": 152323,
+  "<647>": 152324,
+  "<648>": 152325,
+  "<649>": 152326,
+  "<64>": 151741,
+  "<650>": 152327,
+  "<651>": 152328,
+  "<652>": 152329,
+  "<653>": 152330,
+  "<654>": 152331,
+  "<655>": 152332,
+  "<656>": 152333,
+  "<657>": 152334,
+  "<658>": 152335,
+  "<659>": 152336,
+  "<65>": 151742,
+  "<660>": 152337,
+  "<661>": 152338,
+  "<662>": 152339,
+  "<663>": 152340,
+  "<664>": 152341,
+  "<665>": 152342,
+  "<666>": 152343,
+  "<667>": 152344,
+  "<668>": 152345,
+  "<669>": 152346,
+  "<66>": 151743,
+  "<670>": 152347,
+  "<671>": 152348,
+  "<672>": 152349,
+  "<673>": 152350,
+  "<674>": 152351,
+  "<675>": 152352,
+  "<676>": 152353,
+  "<677>": 152354,
+  "<678>": 152355,
+  "<679>": 152356,
+  "<67>": 151744,
+  "<680>": 152357,
+  "<681>": 152358,
+  "<682>": 152359,
+  "<683>": 152360,
+  "<684>": 152361,
+  "<685>": 152362,
+  "<686>": 152363,
+  "<687>": 152364,
+  "<688>": 152365,
+  "<689>": 152366,
+  "<68>": 151745,
+  "<690>": 152367,
+  "<691>": 152368,
+  "<692>": 152369,
+  "<693>": 152370,
+  "<694>": 152371,
+  "<695>": 152372,
+  "<696>": 152373,
+  "<697>": 152374,
+  "<698>": 152375,
+  "<699>": 152376,
+  "<69>": 151746,
+  "<6>": 151683,
+  "<700>": 152377,
+  "<701>": 152378,
+  "<702>": 152379,
+  "<703>": 152380,
+  "<704>": 152381,
+  "<705>": 152382,
+  "<706>": 152383,
+  "<707>": 152384,
+  "<708>": 152385,
+  "<709>": 152386,
+  "<70>": 151747,
+  "<710>": 152387,
+  "<711>": 152388,
+  "<712>": 152389,
+  "<713>": 152390,
+  "<714>": 152391,
+  "<715>": 152392,
+  "<716>": 152393,
+  "<717>": 152394,
+  "<718>": 152395,
+  "<719>": 152396,
+  "<71>": 151748,
+  "<720>": 152397,
+  "<721>": 152398,
+  "<722>": 152399,
+  "<723>": 152400,
+  "<724>": 152401,
+  "<725>": 152402,
+  "<726>": 152403,
+  "<727>": 152404,
+  "<728>": 152405,
+  "<729>": 152406,
+  "<72>": 151749,
+  "<730>": 152407,
+  "<731>": 152408,
+  "<732>": 152409,
+  "<733>": 152410,
+  "<734>": 152411,
+  "<735>": 152412,
+  "<736>": 152413,
+  "<737>": 152414,
+  "<738>": 152415,
+  "<739>": 152416,
+  "<73>": 151750,
+  "<740>": 152417,
+  "<741>": 152418,
+  "<742>": 152419,
+  "<743>": 152420,
+  "<744>": 152421,
+  "<745>": 152422,
+  "<746>": 152423,
+  "<747>": 152424,
+  "<748>": 152425,
+  "<749>": 152426,
+  "<74>": 151751,
+  "<750>": 152427,
+  "<751>": 152428,
+  "<752>": 152429,
+  "<753>": 152430,
+  "<754>": 152431,
+  "<755>": 152432,
+  "<756>": 152433,
+  "<757>": 152434,
+  "<758>": 152435,
+  "<759>": 152436,
+  "<75>": 151752,
+  "<760>": 152437,
+  "<761>": 152438,
+  "<762>": 152439,
+  "<763>": 152440,
+  "<764>": 152441,
+  "<765>": 152442,
+  "<766>": 152443,
+  "<767>": 152444,
+  "<768>": 152445,
+  "<769>": 152446,
+  "<76>": 151753,
+  "<770>": 152447,
+  "<771>": 152448,
+  "<772>": 152449,
+  "<773>": 152450,
+  "<774>": 152451,
+  "<775>": 152452,
+  "<776>": 152453,
+  "<777>": 152454,
+  "<778>": 152455,
+  "<779>": 152456,
+  "<77>": 151754,
+  "<780>": 152457,
+  "<781>": 152458,
+  "<782>": 152459,
+  "<783>": 152460,
+  "<784>": 152461,
+  "<785>": 152462,
+  "<786>": 152463,
+  "<787>": 152464,
+  "<788>": 152465,
+  "<789>": 152466,
+  "<78>": 151755,
+  "<790>": 152467,
+  "<791>": 152468,
+  "<792>": 152469,
+  "<793>": 152470,
+  "<794>": 152471,
+  "<795>": 152472,
+  "<796>": 152473,
+  "<797>": 152474,
+  "<798>": 152475,
+  "<799>": 152476,
+  "<79>": 151756,
+  "<7>": 151684,
+  "<800>": 152477,
+  "<801>": 152478,
+  "<802>": 152479,
+  "<803>": 152480,
+  "<804>": 152481,
+  "<805>": 152482,
+  "<806>": 152483,
+  "<807>": 152484,
+  "<808>": 152485,
+  "<809>": 152486,
+  "<80>": 151757,
+  "<810>": 152487,
+  "<811>": 152488,
+  "<812>": 152489,
+  "<813>": 152490,
+  "<814>": 152491,
+  "<815>": 152492,
+  "<816>": 152493,
+  "<817>": 152494,
+  "<818>": 152495,
+  "<819>": 152496,
+  "<81>": 151758,
+  "<820>": 152497,
+  "<821>": 152498,
+  "<822>": 152499,
+  "<823>": 152500,
+  "<824>": 152501,
+  "<825>": 152502,
+  "<826>": 152503,
+  "<827>": 152504,
+  "<828>": 152505,
+  "<829>": 152506,
+  "<82>": 151759,
+  "<830>": 152507,
+  "<831>": 152508,
+  "<832>": 152509,
+  "<833>": 152510,
+  "<834>": 152511,
+  "<835>": 152512,
+  "<836>": 152513,
+  "<837>": 152514,
+  "<838>": 152515,
+  "<839>": 152516,
+  "<83>": 151760,
+  "<840>": 152517,
+  "<841>": 152518,
+  "<842>": 152519,
+  "<843>": 152520,
+  "<844>": 152521,
+  "<845>": 152522,
+  "<846>": 152523,
+  "<847>": 152524,
+  "<848>": 152525,
+  "<849>": 152526,
+  "<84>": 151761,
+  "<850>": 152527,
+  "<851>": 152528,
+  "<852>": 152529,
+  "<853>": 152530,
+  "<854>": 152531,
+  "<855>": 152532,
+  "<856>": 152533,
+  "<857>": 152534,
+  "<858>": 152535,
+  "<859>": 152536,
+  "<85>": 151762,
+  "<860>": 152537,
+  "<861>": 152538,
+  "<862>": 152539,
+  "<863>": 152540,
+  "<864>": 152541,
+  "<865>": 152542,
+  "<866>": 152543,
+  "<867>": 152544,
+  "<868>": 152545,
+  "<869>": 152546,
+  "<86>": 151763,
+  "<870>": 152547,
+  "<871>": 152548,
+  "<872>": 152549,
+  "<873>": 152550,
+  "<874>": 152551,
+  "<875>": 152552,
+  "<876>": 152553,
+  "<877>": 152554,
+  "<878>": 152555,
+  "<879>": 152556,
+  "<87>": 151764,
+  "<880>": 152557,
+  "<881>": 152558,
+  "<882>": 152559,
+  "<883>": 152560,
+  "<884>": 152561,
+  "<885>": 152562,
+  "<886>": 152563,
+  "<887>": 152564,
+  "<888>": 152565,
+  "<889>": 152566,
+  "<88>": 151765,
+  "<890>": 152567,
+  "<891>": 152568,
+  "<892>": 152569,
+  "<893>": 152570,
+  "<894>": 152571,
+  "<895>": 152572,
+  "<896>": 152573,
+  "<897>": 152574,
+  "<898>": 152575,
+  "<899>": 152576,
+  "<89>": 151766,
+  "<8>": 151685,
+  "<900>": 152577,
+  "<901>": 152578,
+  "<902>": 152579,
+  "<903>": 152580,
+  "<904>": 152581,
+  "<905>": 152582,
+  "<906>": 152583,
+  "<907>": 152584,
+  "<908>": 152585,
+  "<909>": 152586,
+  "<90>": 151767,
+  "<910>": 152587,
+  "<911>": 152588,
+  "<912>": 152589,
+  "<913>": 152590,
+  "<914>": 152591,
+  "<915>": 152592,
+  "<916>": 152593,
+  "<917>": 152594,
+  "<918>": 152595,
+  "<919>": 152596,
+  "<91>": 151768,
+  "<920>": 152597,
+  "<921>": 152598,
+  "<922>": 152599,
+  "<923>": 152600,
+  "<924>": 152601,
+  "<925>": 152602,
+  "<926>": 152603,
+  "<927>": 152604,
+  "<928>": 152605,
+  "<929>": 152606,
+  "<92>": 151769,
+  "<930>": 152607,
+  "<931>": 152608,
+  "<932>": 152609,
+  "<933>": 152610,
+  "<934>": 152611,
+  "<935>": 152612,
+  "<936>": 152613,
+  "<937>": 152614,
+  "<938>": 152615,
+  "<939>": 152616,
+  "<93>": 151770,
+  "<940>": 152617,
+  "<941>": 152618,
+  "<942>": 152619,
+  "<943>": 152620,
+  "<944>": 152621,
+  "<945>": 152622,
+  "<946>": 152623,
+  "<947>": 152624,
+  "<948>": 152625,
+  "<949>": 152626,
+  "<94>": 151771,
+  "<950>": 152627,
+  "<951>": 152628,
+  "<952>": 152629,
+  "<953>": 152630,
+  "<954>": 152631,
+  "<955>": 152632,
+  "<956>": 152633,
+  "<957>": 152634,
+  "<958>": 152635,
+  "<959>": 152636,
+  "<95>": 151772,
+  "<960>": 152637,
+  "<961>": 152638,
+  "<962>": 152639,
+  "<963>": 152640,
+  "<964>": 152641,
+  "<965>": 152642,
+  "<966>": 152643,
+  "<967>": 152644,
+  "<968>": 152645,
+  "<969>": 152646,
+  "<96>": 151773,
+  "<970>": 152647,
+  "<971>": 152648,
+  "<972>": 152649,
+  "<973>": 152650,
+  "<974>": 152651,
+  "<975>": 152652,
+  "<976>": 152653,
+  "<977>": 152654,
+  "<978>": 152655,
+  "<979>": 152656,
+  "<97>": 151774,
+  "<980>": 152657,
+  "<981>": 152658,
+  "<982>": 152659,
+  "<983>": 152660,
+  "<984>": 152661,
+  "<985>": 152662,
+  "<986>": 152663,
+  "<987>": 152664,
+  "<988>": 152665,
+  "<989>": 152666,
+  "<98>": 151775,
+  "<990>": 152667,
+  "<991>": 152668,
+  "<992>": 152669,
+  "<993>": 152670,
+  "<994>": 152671,
+  "<995>": 152672,
+  "<996>": 152673,
+  "<997>": 152674,
+  "<998>": 152675,
+  "<999>": 152676,
+  "<99>": 151776,
+  "<9>": 151686,
+  "<IMG_CONTEXT>": 151665,
+  "<box>": 151668,
+  "<img>": 151666,
+  "<interval>": 151674,
+  "<null>": 152678,
+  "<quad>": 151670,
+  "<ref>": 151672,
+  "<switch>": 152679,
+  "<text_mask>": 151676,
+  "<tool_call>": 151657,
+  "<|box_end|>": 151649,
+  "<|box_start|>": 151648,
+  "<|endoftext|>": 151643,
+  "<|file_sep|>": 151664,
+  "<|fim_middle|>": 151660,
+  "<|fim_pad|>": 151662,
+  "<|fim_prefix|>": 151659,
+  "<|fim_suffix|>": 151661,
+  "<|im_end|>": 151645,
+  "<|im_start|>": 151644,
+  "<|image_pad|>": 151655,
+  "<|object_ref_end|>": 151647,
+  "<|object_ref_start|>": 151646,
+  "<|quad_end|>": 151651,
+  "<|quad_start|>": 151650,
+  "<|repo_name|>": 151663,
+  "<|video_pad|>": 151656,
+  "<|vision_end|>": 151653,
+  "<|vision_pad|>": 151654,
+  "<|vision_start|>": 151652
+}

all_results.json

@@ -0,0 +1,9 @@
+{
+    "epoch": 0.069,
+    "total_flos": 5.6478980657499236e+20,
+    "train_loss": 0.02277738807797432,
+    "train_runtime": 3535.7374,
+    "train_samples": "streaming",
+    "train_samples_per_second": 362.018,
+    "train_steps_per_second": 1.414
+}

assets/decoding_demo.mp4

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

assets/demo.mp4

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:56e5c216d26741be466562e90bf6f3d576f6e3c7ac64341a9101cccf94af2335
+size 101601160

batch_infer.py

@@ -0,0 +1,129 @@
+#!/usr/bin/env python3
+"""Minimal batch inference CLI for the LocateAnything-3B release code.
+
+Examples:
+  python batch_infer.py --model /path/to/LocateAnything-3B --attn sdpa \
+    --image demo.jpg --query "person</c>car"
+
+  python batch_infer.py --requests requests.jsonl --batch-size 16 --attn la_flash
+
+Each JSONL request should contain {"image": "/path/to.jpg", "query": "person</c>car"}.
+"""
+import argparse
+import json
+import os
+from pathlib import Path
+
+from PIL import Image
+
+
+def _attn_arg(value):
+    mode = (value or "sdpa").strip().lower().replace("-", "_")
+    aliases = {
+        "": "sdpa",
+        "manual": "eager",
+        "torch": "eager",
+        "torch_eager": "eager",
+        "torch_sdpa": "sdpa",
+        "flash": "la_flash",
+        "la_flash": "la_flash",
+        "kernel": "la_flash",
+        "cuda": "la_flash",
+        "range": "la_flash",
+        "range_attention": "la_flash",
+    }
+    mode = aliases.get(mode, mode)
+    if mode not in {"sdpa", "eager", "magi", "la_flash"}:
+        raise argparse.ArgumentTypeError(
+            f"--attn must be one of sdpa, eager, magi, la_flash; got {value!r}"
+        )
+    return mode
+
+
+def _load_requests(args):
+    requests = []
+    if args.requests:
+        with open(args.requests, "r", encoding="utf-8") as f:
+            for line in f:
+                if not line.strip():
+                    continue
+                row = json.loads(line)
+                requests.append((row["image"], row["query"]))
+    if args.image or args.query:
+        if len(args.image or []) != len(args.query or []):
+            raise ValueError("--image and --query must appear the same number of times")
+        requests.extend(zip(args.image, args.query))
+    if not requests:
+        raise ValueError("provide --requests JSONL or at least one --image/--query pair")
+    return requests
+
+
+def main():
+    ap = argparse.ArgumentParser()
+    ap.add_argument("--requests", help="JSONL file with image/query fields")
+    ap.add_argument("--image", action="append", help="Image path; repeat with --query")
+    ap.add_argument("--query", action="append", help="Category query, e.g. person</c>car")
+    ap.add_argument("--model", default=os.environ.get("LA_FLASH_MODEL", "nvidia/LocateAnything-3B"))
+    ap.add_argument("--attn", type=_attn_arg, default=os.environ.get("LA_FLASH_ATTN", "sdpa"),
+                    help="LLM attention backend: sdpa, eager, magi, or la_flash")
+    ap.add_argument("--vision-attn", default=os.environ.get("LA_FLASH_VISION_ATTN", "auto"),
+                    choices=["auto", "flash_attention_2", "sdpa", "eager"])
+    ap.add_argument("--batch-size", type=int, default=1)
+    ap.add_argument("--scheduler", default=os.environ.get("LA_FLASH_HYBRID_SCHEDULER", "eager"),
+                    choices=["eager", "hold_ar", "ar_first", "pipeline", "adaptive"])
+    ap.add_argument("--group-size", type=int, default=int(os.environ.get("LA_FLASH_HYBRID_GROUP_SIZE", "0")))
+    ap.add_argument("--max-new-tokens", type=int, default=2048)
+    ap.add_argument("--temperature", type=float, default=0.7)
+    ap.add_argument("--top-p", type=float, default=0.9)
+    ap.add_argument("--top-k", type=int, default=0)
+    ap.add_argument("--repetition-penalty", type=float, default=1.1)
+    ap.add_argument("--strict-attn", action="store_true",
+                    help="Fail instead of falling back to SDPA if magi/la_flash is unavailable")
+    ap.add_argument("--out", default="", help="Optional output JSONL path; stdout if omitted")
+    args = ap.parse_args()
+    args.attn = _attn_arg(args.attn)
+
+    os.environ["LA_FLASH_MODEL"] = args.model
+    os.environ["LA_FLASH_ATTN"] = args.attn
+    os.environ["LA_FLASH_VISION_ATTN"] = args.vision_attn
+    os.environ["LA_FLASH_HYBRID_SCHEDULER"] = args.scheduler
+    os.environ["LA_FLASH_HYBRID_GROUP_SIZE"] = str(args.group_size)
+    if args.strict_attn:
+        os.environ["LA_FLASH_STRICT_ATTN"] = "1"
+
+    from batch_utils import generate_batch_hybrid, get_last_hybrid_stats, load
+    from batch_utils.hybrid_runtime import load_pil
+
+    requests = _load_requests(args)
+    load()
+
+    writer = open(args.out, "w", encoding="utf-8") if args.out else None
+    try:
+        for start in range(0, len(requests), max(1, args.batch_size)):
+            chunk = requests[start:start + max(1, args.batch_size)]
+            pairs = [(load_pil(image), query) for image, query in chunk]
+            texts = generate_batch_hybrid(
+                pairs,
+                temperature=args.temperature,
+                top_p=None if args.top_p < 0 else args.top_p,
+                top_k=None if args.top_k <= 0 else args.top_k,
+                repetition_penalty=args.repetition_penalty,
+                max_new_tokens=args.max_new_tokens,
+                scheduler=args.scheduler,
+                group_size=args.group_size,
+            )
+            stats = get_last_hybrid_stats()
+            for (image, query), text in zip(chunk, texts):
+                row = {"image": str(Path(image)), "query": query, "raw_response": text, "stats": stats}
+                line = json.dumps(row, ensure_ascii=False)
+                if writer:
+                    writer.write(line + "\n")
+                else:
+                    print(line, flush=True)
+    finally:
+        if writer:
+            writer.close()
+
+
+if __name__ == "__main__":
+    main()

batch_utils/README.md

@@ -0,0 +1,53 @@
+# Batch Utils
+
+`batch_utils` contains the optional batched hybrid generation runtime for
+LocateAnything. It keeps the model loading, tokenization, image feature caching,
+sampling, and scheduler code used by `batch_infer.py` and the detection
+experiments.
+
+## Runtime Modes
+
+- `LA_FLASH_ATTN=sdpa`: stock PyTorch SDPA path.
+- `LA_FLASH_ATTN=eager`: eager attention path for debugging.
+- `LA_FLASH_ATTN=magi`: MagiAttention path when MagiAttention is installed.
+- `LA_FLASH_ATTN=la_flash`: LA Flash sparse range backend
+  from `kernel_utils`.
+
+## Common Knobs
+
+| Variable | Default | Meaning |
+| --- | --- | --- |
+| `LA_FLASH_MODEL` | `nvidia/LocateAnything-3B` | HF model id or local model directory. |
+| `LA_FLASH_ATTN` | `sdpa` | LLM attention backend. |
+| `LA_FLASH_VISION_ATTN` | `auto` | Vision encoder attention: `auto`, `flash_attention_2`, `sdpa`, or `eager`. |
+| `LA_FLASH_STRICT_ATTN` | `0` | Set `1` to fail instead of falling back to SDPA. |
+| `LA_FLASH_HYBRID_SCHEDULER` | `eager` | Hybrid decode scheduler. |
+| `LA_FLASH_HYBRID_GROUP_SIZE` | `0` | Scheduler group size; `0` lets the runtime decide. |
+| `LA_FLASH_VISION_ENCODE_BATCH_SIZE` | `8` | Maximum images per MoonViT encode micro-batch. |
+| `LA_FLASH_KV_PACK_TOKEN_BUDGET` | `0` | Optional KV packing memory cap for long-tail batches. |
+| `LA_FLASH_DENSE_BACKEND` | `sdpa` | Dense worker/prefill attention backend. Keep this as `sdpa`; LA Flash is used for sparse range plans. |
+| `LA_FLASH_SEGMENT_FASTPATH` | `auto` | Sparse MTP decode uses FlashAttention varlen multi-segment merge by default. |
+
+## CLI Example
+
+```bash
+python batch_infer.py \
+  --model nvidia/LocateAnything-3B \
+  --attn la_flash \
+  --scheduler pipeline \
+  --batch-size 4 \
+  --image /path/to/image.jpg \
+  --query "person</c>car"
+```
+
+For JSONL input, each row should contain:
+
+```json
+{"image": "/path/to/image.jpg", "query": "person</c>car"}
+```
+
+## Training Boundary
+
+This package is for inference and evaluation. Training remains on the
+MagiAttention backend; the batched sparse-plan decode runtime does not support
+the `labels` training path.

batch_utils/__init__.py

@@ -0,0 +1,12 @@
+"""Hybrid batched inference helpers for NVIDIA LocateAnything-3B."""
+from .hybrid_runtime import load, load_pil
+from .engine_hybrid import generate_batch_hybrid, generate_batch_grouped_hybrid, get_last_hybrid_stats
+
+__all__ = [
+    "load",
+    "load_pil",
+    "generate_batch_hybrid",
+    "generate_batch_grouped_hybrid",
+    "get_last_hybrid_stats",
+]
+__version__ = "0.1.0"

batch_utils/engine_hybrid.py

@@ -0,0 +1,1357 @@
+"""Batched hybrid-mode generation for LocateAnything-3B.
+
+This module keeps the stock hybrid state machine:
+
+  MTP -> error_box -> AR
+  AR  -> box_end_ar -> MTP
+
+Rows in a batch may be in different modes. The decode loop therefore stores
+per-row KV caches, packs rows with the same mode for one forward call, then
+unpacks the clean KV back per row.
+"""
+import copy
+import importlib
+import os
+
+import torch
+
+
+from .hybrid_runtime import (
+    ATTN_MODE,
+    AR_BATCH_SAN,
+    BATCH_SAN,
+    DEV,
+    N_FUTURE,
+    _encode_images,
+    _helpers,
+    _pad_generated,
+    _set_llm_mode,
+    _tokenize,
+    _tokenize_cached_image,
+    build_magi_scheduler_ranges,
+    language_model_forward,
+    load,
+    sample_next_tokens_batched,
+    sample_tokens_batched,
+)
+
+
+README_MAX_NEW_TOKENS = 2048
+README_TEMPERATURE = 0.7
+README_TOP_P = 0.9
+README_REPETITION_PENALTY = 1.1
+
+_LAST_HYBRID_STATS = None
+
+
+def _row_len(kv):
+    return kv[0][0].shape[2]
+
+
+def _pack_stock_kv_rows(kv_rows, rows, dev):
+    """Left-pad per-row real-token KV caches for stock-style decoding."""
+    lengths = [0 if kv_rows[r] is None else _row_len(kv_rows[r]) for r in rows]
+    kmax = max(lengths) if lengths else 0
+    if kmax == 0:
+        return None, torch.zeros((len(rows), 0), dtype=torch.long, device=dev), lengths, 0
+
+    ref = next(kv_rows[r] for r in rows if kv_rows[r] is not None)
+    packed = []
+    for layer in range(len(ref)):
+        ref_k, ref_v = ref[layer]
+        ks, vs = [], []
+        for r, length in zip(rows, lengths):
+            if length == 0:
+                k = ref_k.new_zeros((1, ref_k.shape[1], kmax, ref_k.shape[3]))
+                v = ref_v.new_zeros((1, ref_v.shape[1], kmax, ref_v.shape[3]))
+            else:
+                k, v = kv_rows[r][layer]
+                if length < kmax:
+                    pad_shape = (1, k.shape[1], kmax - length, k.shape[3])
+                    k = torch.cat([k.new_zeros(pad_shape), k], dim=2)
+                    v = torch.cat([v.new_zeros(pad_shape), v], dim=2)
+            ks.append(k)
+            vs.append(v)
+        packed.append((torch.cat(ks, dim=0), torch.cat(vs, dim=0)))
+
+    kvalid = torch.zeros((len(rows), kmax), dtype=torch.long, device=dev)
+    for i, length in enumerate(lengths):
+        if length:
+            kvalid[i, kmax - length :] = 1
+    return tuple(packed), kvalid, lengths, kmax
+
+
+def _unpack_stock_after_forward(out_kv, local_row, old_len, uncached_len, kmax, umax):
+    """Keep old real KV plus the right-aligned uncached real tokens; drop pads/window."""
+    out = []
+    u0 = kmax + (umax - uncached_len)
+    u1 = kmax + umax
+    for k, v in out_kv:
+        parts_k, parts_v = [], []
+        if old_len:
+            parts_k.append(k[local_row : local_row + 1, :, kmax - old_len : kmax, :])
+            parts_v.append(v[local_row : local_row + 1, :, kmax - old_len : kmax, :])
+        if uncached_len:
+            parts_k.append(k[local_row : local_row + 1, :, u0:u1, :])
+            parts_v.append(v[local_row : local_row + 1, :, u0:u1, :])
+        out.append((torch.cat(parts_k, dim=2).contiguous(),
+                    torch.cat(parts_v, dim=2).contiguous()))
+    return tuple(out)
+
+
+def _mk_generate_kwargs(temperature, top_p, top_k, repetition_penalty, row_temp=None):
+    t = temperature if row_temp is None else row_temp
+    gk = {"repetition_penalty": repetition_penalty, "generation_mode": "hybrid"}
+    if t and t > 0:
+        gk["temperature"] = t
+    if top_p is not None:
+        gk["top_p"] = top_p
+    if top_k is not None:
+        gk["top_k"] = top_k
+    return gk
+
+
+def _classify_ar_token(token_val, tids):
+    if token_val == tids["box_end_token_id"]:
+        return "box_end_ar"
+    if tids["coord_start_token_id"] <= token_val <= tids["coord_end_token_id"]:
+        return "coord_ar"
+    if token_val == tids["none_token_id"]:
+        return "coord_ar"
+    return "im_end"
+
+
+def _env_flag(name, default=False):
+    val = os.environ.get(name)
+    if val is None:
+        return default
+    return val.lower() not in {"0", "false", "no", "off", ""}
+
+
+def _env_int(name, default):
+    val = os.environ.get(name)
+    if val is None or val == "":
+        return default
+    return int(val)
+
+
+def _kv_pack_token_budget():
+    return max(0, _env_int("LA_FLASH_KV_PACK_TOKEN_BUDGET", 0))
+
+
+def _debug_enabled(debug):
+    return _env_flag("LA_FLASH_DEBUG", False) if debug is None else bool(debug)
+
+
+def _new_hybrid_stats(total_rows, scheduler, group_size, hold_max_steps, adaptive_hold_mtp_max=0):
+    return {
+        "scheduler": scheduler,
+        "requested_group_size": int(group_size or 0),
+        "hold_max_steps": int(hold_max_steps),
+        "adaptive_hold_mtp_max": int(adaptive_hold_mtp_max),
+        "input_batches": 1,
+        "input_rows": int(total_rows),
+        "groups": 0,
+        "group_sizes": [],
+        "decode_loops": 0,
+        "mixed_mode_cycles": 0,
+        "eager_mtp_then_ar_cycles": 0,
+        "ar_first_cycles": 0,
+        "pipeline_ar_after_mtp_cycles": 0,
+        "adaptive_hold_cycles": 0,
+        "adaptive_ar_first_cycles": 0,
+        "hold_ar_steps": 0,
+        "hold_ar_held_mtp_rows": 0,
+        "hold_ar_limit_mtp_forwards": 0,
+        "mtp_forwards": 0,
+        "ar_forwards": 0,
+        "mtp_forward_rows": 0,
+        "ar_forward_rows": 0,
+        "mtp_forward_query_tokens": 0,
+        "ar_forward_query_tokens": 0,
+        "max_mtp_forward_rows": 0,
+        "max_ar_forward_rows": 0,
+        "mtp_max_uncached_len": 0,
+        "ar_max_uncached_len": 0,
+        "mtp_forward_row_hist": {},
+        "ar_forward_row_hist": {},
+        "prompt_prefill_mode": _hybrid_prefill_mode(),
+        "prompt_prefill_forwards": 0,
+        "prompt_prefill_forward_rows": 0,
+        "prompt_prefill_forward_query_tokens": 0,
+        "prompt_prefill_real_tokens": 0,
+        "prompt_prefill_shared_groups": 0,
+        "prompt_prefill_shared_rows": 0,
+        "prompt_prefill_shared_saved_tokens": 0,
+        "kv_bucket_splits": 0,
+        "kv_bucket_groups": 0,
+        "kv_bucket_max_packed_tokens": 0,
+    }
+
+
+def _set_last_hybrid_stats(stats):
+    global _LAST_HYBRID_STATS
+    _LAST_HYBRID_STATS = copy.deepcopy(stats) if stats is not None else None
+
+
+def get_last_hybrid_stats():
+    """Return scheduler/forward statistics from the most recent hybrid batch."""
+    return copy.deepcopy(_LAST_HYBRID_STATS)
+
+
+def _record_group_stats(stats, bsz):
+    if stats is None:
+        return
+    stats["groups"] += 1
+    stats["group_sizes"].append(int(bsz))
+
+
+def _bump_hist(hist, val):
+    key = str(int(val))
+    hist[key] = int(hist.get(key, 0)) + 1
+
+
+def _record_forward_stats(stats, kind, rows, q_len, uncached_lens):
+    if stats is None:
+        return
+    prefix = "mtp" if kind == "mtp" else "ar"
+    nrows = int(len(rows))
+    q_len = int(q_len)
+    stats[f"{prefix}_forwards"] += 1
+    stats[f"{prefix}_forward_rows"] += nrows
+    stats[f"{prefix}_forward_query_tokens"] += nrows * q_len
+    stats[f"max_{prefix}_forward_rows"] = max(stats[f"max_{prefix}_forward_rows"], nrows)
+    stats[f"{prefix}_max_uncached_len"] = max(
+        stats[f"{prefix}_max_uncached_len"],
+        max((int(x) for x in uncached_lens), default=0),
+    )
+    _bump_hist(stats[f"{prefix}_forward_row_hist"], nrows)
+
+
+def _record_prefill_stats(stats, rows, q_len, real_tokens, shared_groups=0, shared_rows=0, saved_tokens=0):
+    if stats is None:
+        return
+    nrows = int(rows)
+    stats["prompt_prefill_forwards"] += 1
+    stats["prompt_prefill_forward_rows"] += nrows
+    stats["prompt_prefill_forward_query_tokens"] += nrows * int(q_len)
+    stats["prompt_prefill_real_tokens"] += int(real_tokens)
+    stats["prompt_prefill_shared_groups"] += int(shared_groups)
+    stats["prompt_prefill_shared_rows"] += int(shared_rows)
+    stats["prompt_prefill_shared_saved_tokens"] += int(saved_tokens)
+
+
+def _split_rows_by_kv_budget(rows, kv_rows):
+    """Keep dense left-padded KV packs bounded when a few rows become long tails."""
+    budget = _kv_pack_token_budget()
+    if budget <= 0 or len(rows) <= 1:
+        return [rows]
+    lengths = [0 if kv_rows[r] is None else _row_len(kv_rows[r]) for r in rows]
+    if not lengths or max(lengths) * len(rows) <= budget:
+        return [rows]
+
+    groups = []
+    current = []
+    current_max = 0
+    for row, length in sorted(zip(rows, lengths), key=lambda item: item[1]):
+        next_max = max(current_max, int(length))
+        if current and next_max * (len(current) + 1) > budget:
+            groups.append(current)
+            current = [row]
+            current_max = int(length)
+        else:
+            current.append(row)
+            current_max = next_max
+    if current:
+        groups.append(current)
+    return groups or [rows]
+
+
+def _record_kv_bucket_stats(stats, groups, kv_rows):
+    if stats is None:
+        return
+    max_packed = 0
+    for group in groups:
+        if not group:
+            continue
+        kmax = max((0 if kv_rows[r] is None else _row_len(kv_rows[r])) for r in group)
+        max_packed = max(max_packed, int(kmax) * len(group))
+    stats["kv_bucket_max_packed_tokens"] = max(stats["kv_bucket_max_packed_tokens"], max_packed)
+    if len(groups) > 1:
+        stats["kv_bucket_splits"] += 1
+        stats["kv_bucket_groups"] += len(groups)
+
+
+def _hybrid_scheduler(scheduler):
+    val = os.environ.get("LA_FLASH_HYBRID_SCHEDULER", "eager") if scheduler is None else scheduler
+    val = str(val).strip().lower()
+    aliases = {
+        "": "eager",
+        "default": "eager",
+        "normal": "eager",
+        "hold": "hold_ar",
+        "hold-ar": "hold_ar",
+        "hold_mtp": "hold_ar",
+        "hold-mtp": "hold_ar",
+        "repair_first": "ar_first",
+        "repair-first": "ar_first",
+        "ar-first": "ar_first",
+    }
+    val = aliases.get(val, val)
+    if val not in {"eager", "hold_ar", "ar_first", "pipeline", "adaptive"}:
+        raise ValueError("scheduler must be one of: eager, hold_ar, ar_first, pipeline, adaptive")
+    return val
+
+
+def _hybrid_group_size(group_size):
+    if group_size is None:
+        return max(0, _env_int("LA_FLASH_HYBRID_GROUP_SIZE", 0))
+    return max(0, int(group_size))
+
+
+def _hybrid_prefill_mode():
+    val = os.environ.get("LA_FLASH_HYBRID_PREFILL", "shared").strip().lower()
+    aliases = {
+        "0": "none",
+        "false": "none",
+        "off": "none",
+        "legacy": "none",
+        "1": "per_row",
+        "true": "per_row",
+        "on": "per_row",
+        "single": "per_row",
+        "row": "per_row",
+        "rows": "per_row",
+        "batched": "batch",
+        "prefix": "shared",
+        "shared_prefix": "shared",
+        "shared-image": "shared",
+        "shared_image": "shared",
+        "vision": "shared",
+    }
+    val = aliases.get(val, val)
+    if val not in {"none", "per_row", "batch", "shared"}:
+        raise ValueError("LA_FLASH_HYBRID_PREFILL must be one of none, per_row, batch, shared")
+    return val
+
+
+def _tolist(t):
+    return t.detach().cpu().tolist()
+
+
+def _safe_decode_rows(tok, input_ids):
+    rows = []
+    for row in _tolist(input_ids):
+        try:
+            rows.append(tok.decode(torch.tensor(row), skip_special_tokens=False))
+        except Exception:
+            rows.append("<decode failed>")
+    return rows
+
+
+def _safe_decode_row(tok, row):
+    try:
+        return tok.decode(torch.tensor(row), skip_special_tokens=False)
+    except Exception:
+        return "<decode failed>"
+
+
+def _effective_allowed_mask(mask2d, q_len, past_len, mtp_window=False):
+    """Readable 1/0 q-by-k mask derived from the 2D key-valid mask.
+
+    This mirrors the model path at a high level:
+    causal + padding columns, then the MTP window update
+    attn[-block:, -block:] = visible and attn[-block:, -block-1] = masked.
+    """
+    rows = []
+    key_valid = mask2d.detach().cpu().bool()
+    total_len = int(key_valid.numel())
+    for qi in range(q_len):
+        q_abs = past_len + qi
+        row = []
+        for ki in range(total_len):
+            row.append(1 if bool(key_valid[ki]) and ki <= q_abs else 0)
+        rows.append(row)
+
+    if mtp_window and q_len >= N_FUTURE and total_len >= N_FUTURE:
+        q0 = q_len - N_FUTURE
+        k0 = total_len - N_FUTURE
+        for qi in range(q0, q_len):
+            for ki in range(k0, total_len):
+                rows[qi][ki] = 1
+            if k0 - 1 >= 0:
+                rows[qi][k0 - 1] = 0
+    return rows
+
+
+def _tail_matrix(mat, rows=None, cols=None):
+    if rows is not None:
+        mat = mat[-rows:]
+    if cols is not None:
+        mat = [row[-cols:] for row in mat]
+    return mat
+
+
+def _format_01_matrix(mat):
+    return "\n".join("  " + " ".join(str(int(v)) for v in row) for row in mat)
+
+
+def _safe_sdpa_mask_enabled():
+    return _env_flag("LA_FLASH_SDPA_SAFE_4D_MASK", True)
+
+
+def _build_safe_sdpa_visible_mask(attention_mask_2d, input_ids, past_len, mtp_window=False):
+    """Build a 4D 1/0 visible mask, with harmless visibility for all-masked pad queries.
+
+    The remote Qwen2 SDPA path uses a 2D key-valid mask and can create fully
+    masked query rows for left-padded, no-cache prefill. Those rows can produce
+    NaNs inside SDPA and later contaminate real tokens through masked K columns.
+    This 4D mask keeps real-token visibility identical, and only gives otherwise
+    all-masked query rows one valid fallback key so their activations stay finite.
+    """
+    bsz, q_len = int(input_ids.shape[0]), int(input_ids.shape[1])
+    key_len = int(attention_mask_2d.shape[1])
+    dev = input_ids.device
+    key_valid = attention_mask_2d.to(dtype=torch.bool, device=dev)
+    key_idx = torch.arange(key_len, device=dev).view(1, 1, key_len)
+    q_abs = (past_len + torch.arange(q_len, device=dev)).view(1, q_len, 1)
+    visible = key_valid[:, None, :] & (key_idx <= q_abs)
+
+    if mtp_window and q_len >= N_FUTURE and key_len >= N_FUTURE:
+        k0 = key_len - N_FUTURE
+        visible[:, -N_FUTURE:, k0:key_len] = key_valid[:, None, k0:key_len]
+        blocked_k = k0 - 1
+        if blocked_k >= 0:
+            visible[:, -N_FUTURE:, blocked_k] = False
+
+    row_has_key = visible.any(dim=-1)
+    fallback_rows = int((~row_has_key).sum().item())
+    if fallback_rows:
+        for b in range(bsz):
+            valid = torch.nonzero(key_valid[b], as_tuple=False).flatten()
+            fallback = int(valid[0].item()) if valid.numel() else 0
+            missing = torch.nonzero(~row_has_key[b], as_tuple=False).flatten()
+            if missing.numel():
+                visible[b, missing, fallback] = True
+
+    mask = visible[:, None, :, :].to(dtype=torch.bfloat16)
+    try:
+        mask._la_flash_visible_mask = True
+    except Exception:
+        pass
+    return mask, fallback_rows
+
+
+def _mask_desc(mask):
+    if mask is None:
+        return "none"
+    if isinstance(mask, dict):
+        return "magi_ranges"
+    if hasattr(mask, "dim"):
+        return "4d_safe_sdpa" if mask.dim() == 4 else "2d_key_valid"
+    return type(mask).__name__
+
+
+def _forward_attention_mask(model, input_ids, attention_mask_2d, past_len, mtp_window=False, range_plan=False):
+    llm = model.language_model.model
+    if getattr(model, "_la_flash_requested_attn", ATTN_MODE) in {"magi", "la_flash"}:
+        range_plan = build_magi_scheduler_ranges(
+            model, attention_mask_2d, input_ids, past_len, mtp_window=mtp_window)
+        if range_plan is not None:
+            return range_plan, 0
+    needs_safe_pad = (
+        past_len == 0
+        and attention_mask_2d is not None
+        and attention_mask_2d.dim() == 2
+        and input_ids.shape[0] > 1
+    )
+    if (
+        getattr(llm, "_attn_implementation", None) == "sdpa"
+        and _safe_sdpa_mask_enabled()
+        and needs_safe_pad
+        and attention_mask_2d is not None
+        and attention_mask_2d.dim() == 2
+    ):
+        return _build_safe_sdpa_visible_mask(attention_mask_2d, input_ids, past_len, mtp_window)
+    return attention_mask_2d, 0
+
+
+def _actual_sdpa_allowed_masks(model, input_ids, attention_mask, past_len):
+    """Recreate the remote Qwen2 SDPA 4D additive mask and return a 0/1 view."""
+    llm = model.language_model.model
+    mod = importlib.import_module(type(llm).__module__)
+    bsz, q_len = int(input_ids.shape[0]), int(input_ids.shape[1])
+    dummy = torch.empty(
+        (bsz, q_len, 1),
+        dtype=torch.bfloat16,
+        device=input_ids.device,
+    )
+    mask4 = mod._prepare_4d_causal_attention_mask(
+        attention_mask,
+        (bsz, q_len),
+        dummy,
+        past_len,
+        sliding_window=getattr(llm.config, "sliding_window", None),
+    )
+    remote_ar_decode = q_len == 1 or (
+        input_ids is not None and int(input_ids[0, -1].item()) != int(llm.text_mask_token_id)
+    )
+    if not remote_ar_decode and mask4 is not None and mask4.dim() == 4:
+        rows = []
+        for b in range(bsz):
+            rows.append(
+                mod.update_causal_mask_for_one_gen_window_2d(
+                    input_ids[b],
+                    mask4[b][0].clone(),
+                    block_size=int(llm.block_size),
+                    use_cache=True,
+                    causal_attn=bool(getattr(llm, "causal_attn", False)),
+                ).unsqueeze(0)
+            )
+        mask4 = torch.stack(rows, dim=0)
+    allowed = (mask4[:, 0] >= 0).to(torch.int8).detach().cpu().tolist()
+    return allowed, tuple(mask4.shape), remote_ar_decode
+
+
+def _debug_magi_ranges(q_len, past_len, mtp_window=False):
+    kv_len = past_len + q_len
+    ar_decode = not mtp_window
+    if ar_decode:
+        return {
+            "q_ranges": [[0, q_len]],
+            "k_ranges": [[0, kv_len]],
+            "attn_type_map": ["CAUSAL"],
+        }
+
+    block = N_FUTURE
+    if not (0 < block <= q_len <= kv_len):
+        return {"error": f"invalid magi MTP shape: block={block}, q_len={q_len}, kv_len={kv_len}"}
+
+    prefix_len = kv_len - block
+    blocked_k = prefix_len - 1
+    q_ranges, k_ranges, attn_types = [], [], []
+    if q_len == kv_len:
+        if prefix_len > 0:
+            q_ranges.append([0, prefix_len])
+            k_ranges.append([0, prefix_len])
+            attn_types.append("CAUSAL")
+        if prefix_len > 0 and blocked_k > 0:
+            q_ranges.append([prefix_len, kv_len])
+            k_ranges.append([0, blocked_k])
+            attn_types.append("FULL")
+        q_ranges.append([prefix_len, kv_len])
+        k_ranges.append([prefix_len, kv_len])
+        attn_types.append("FULL")
+    else:
+        recompute = q_len - block
+        q_global_start = kv_len - q_len
+        for i in range(recompute):
+            g = q_global_start + i
+            q_ranges.append([i, i + 1])
+            k_ranges.append([0, g + 1])
+            attn_types.append("FULL")
+        q_win = [recompute, q_len]
+        if blocked_k > 0:
+            q_ranges.append(q_win)
+            k_ranges.append([0, blocked_k])
+            attn_types.append("FULL")
+        q_ranges.append(q_win)
+        k_ranges.append([prefix_len, kv_len])
+        attn_types.append("FULL")
+
+    return {"q_ranges": q_ranges, "k_ranges": k_ranges, "attn_type_map": attn_types}
+
+
+def _print_debug_forward(label, model, tok, input_ids, attention_mask, position_ids,
+                         past_len, mtp_window=False, extra=None, attention_impl="sdpa"):
+    print(f"\n========== LA Flash DEBUG {label} ==========", flush=True)
+    if extra:
+        for k, v in extra.items():
+            print(f"{k}: {v}", flush=True)
+    tail = int(os.environ.get("LA_FLASH_DEBUG_TAIL", "15"))
+    bsz, q_len = int(input_ids.shape[0]), int(input_ids.shape[1])
+    key_len = int(attention_mask.shape[1])
+    q_tail, k_tail = min(tail, q_len), min(tail, key_len)
+    print(
+        "shapes: "
+        f"input_ids={tuple(input_ids.shape)} "
+        f"position_ids={tuple(position_ids.shape)} "
+        f"attention_mask_key_valid={tuple(attention_mask.shape)} "
+        f"mask_2d_q_by_k=({bsz}, {q_len}, {key_len}) "
+        f"mask_2d_tail=({bsz}, {q_tail}, {k_tail}) "
+        f"past_len={past_len} q_len={q_len} "
+        f"mtp_window={mtp_window} ar_decode={not mtp_window}",
+        flush=True,
+    )
+    print(f"dtypes/devices: input_ids={input_ids.dtype}@{input_ids.device} position_ids={position_ids.dtype}@{position_ids.device} attention_mask={attention_mask.dtype}@{attention_mask.device}", flush=True)
+    print(f"attention_impl={attention_impl}", flush=True)
+    input_rows = _tolist(input_ids)
+    pos_rows = _tolist(position_ids)
+    print(f"tail_window_last={tail}", flush=True)
+    print(f"input_ids_tail.shape=({bsz}, {q_tail})", flush=True)
+    print(f"position_ids_tail.shape=({bsz}, {q_tail})", flush=True)
+    actual_sdpa = None
+    if attention_impl in {"sdpa", "eager", "la_flash"}:
+        try:
+            actual_sdpa = _actual_sdpa_allowed_masks(model, input_ids, attention_mask, past_len)
+            print(
+                f"actual_sdpa_4d_mask_shape={actual_sdpa[1]} "
+                f"remote_ar_decode={actual_sdpa[2]}",
+                flush=True,
+            )
+        except Exception as e:
+            print(f"actual_sdpa_4d_mask_debug_failed={type(e).__name__}: {e}", flush=True)
+
+    for b in range(input_ids.shape[0]):
+        ids_tail = input_rows[b][-tail:]
+        pos_tail = pos_rows[b][-tail:]
+        allowed = _effective_allowed_mask(attention_mask[b], input_ids.shape[1], past_len, mtp_window)
+        q_tail = min(tail, len(allowed))
+        k_tail = min(tail, len(allowed[0]) if allowed else 0)
+        allowed_tail = _tail_matrix(allowed, rows=q_tail, cols=k_tail)
+        print(f"batch_row={b} ar_decode={not mtp_window}", flush=True)
+        print(f"input_ids_tail[-{tail}:]: {ids_tail}", flush=True)
+        print(f"decoded_tail[-{tail}:]: {_safe_decode_row(tok, ids_tail)}", flush=True)
+        print(f"position_ids_tail[-{tail}:]: {pos_tail}", flush=True)
+        print(f"expected_mask_2d_tail[-{q_tail}:,-{k_tail}:].shape=({q_tail}, {k_tail})", flush=True)
+        print(_format_01_matrix(allowed_tail), flush=True)
+        if actual_sdpa is not None:
+            actual = actual_sdpa[0][b]
+            actual_tail = _tail_matrix(actual, rows=q_tail, cols=k_tail)
+            mismatch = sum(
+                int(allowed[qi][ki] != actual[qi][ki])
+                for qi in range(len(allowed))
+                for ki in range(len(allowed[qi]))
+            )
+            print(
+                f"actual_sdpa_mask_2d_tail[-{q_tail}:,-{k_tail}:].shape=({q_tail}, {k_tail})",
+                flush=True,
+            )
+            print(_format_01_matrix(actual_tail), flush=True)
+            print(f"expected_vs_actual_sdpa_mismatch_count={mismatch}", flush=True)
+
+    if _env_flag("LA_FLASH_DEBUG_FULL_MASK", False):
+        masks = [
+            _effective_allowed_mask(attention_mask[b], input_ids.shape[1], past_len, mtp_window)
+            for b in range(input_ids.shape[0])
+        ]
+        print("effective_allowed_mask_q_by_k_FULL:", masks, flush=True)
+    if attention_impl == "magi":
+        if bsz == 1:
+            print(
+                "magi_ranges:",
+                _debug_magi_ranges(input_ids.shape[1], past_len, mtp_window),
+                flush=True,
+            )
+        else:
+            print(
+                "magi_ranges: built once per forward from the batched scheduler mask",
+                flush=True,
+            )
+            print(
+                "magi_ranges_single_row_template:",
+                _debug_magi_ranges(input_ids.shape[1], past_len, mtp_window),
+                flush=True,
+            )
+
+
+def _common_prefix_len(prompt_ids, rows):
+    if not rows:
+        return 0
+    first = prompt_ids[rows[0]]
+    max_len = min(int(prompt_ids[r].numel()) for r in rows)
+    prefix_len = 0
+    for idx in range(max_len):
+        val = int(first[idx].item())
+        if all(int(prompt_ids[r][idx].item()) == val for r in rows[1:]):
+            prefix_len += 1
+        else:
+            break
+    return prefix_len
+
+
+def _prefill_shared_prefix_kv_rows(model, prompt_ids, vit_list, img_tok, pad, dev, stats=None, debug=False):
+    """Cache one common prompt prefix per image-feature group.
+
+    Multi-category split repeats the same image feature tensor for each
+    category prompt.  Token ids are identical through the image tokens and the
+    fixed prompt prefix, so we prefill that shared prefix once and let each
+    category row forward only its text suffix.
+    """
+    bsz = len(prompt_ids)
+    kv_rows = [None] * bsz
+    cached_lens = [0] * bsz
+    groups = {}
+    for row, vit in enumerate(vit_list):
+        groups.setdefault(id(vit), []).append(row)
+
+    items = []
+    min_prefix_len = max(1, _env_int("LA_FLASH_SHARED_PREFILL_MIN_PREFIX", 64))
+    for rows in groups.values():
+        if len(rows) < 2:
+            continue
+        prefix_len = _common_prefix_len(prompt_ids, rows)
+        if prefix_len < min_prefix_len:
+            continue
+        prefix_ids = prompt_ids[rows[0]][:prefix_len]
+        image_token_count = int((prefix_ids == img_tok).sum().item())
+        if image_token_count != int(vit_list[rows[0]].shape[0]):
+            if debug:
+                print(
+                    "LA Flash shared prefill skip group: "
+                    f"rows={rows} prefix_len={prefix_len} "
+                    f"image_tokens={image_token_count} visual_rows={int(vit_list[rows[0]].shape[0])}",
+                    flush=True,
+                )
+            continue
+        items.append((rows, prefix_ids, vit_list[rows[0]]))
+
+    if not items:
+        return kv_rows, cached_lens
+
+    lengths = [int(ids.numel()) for _rows, ids, _vit in items]
+    pmax = max(lengths)
+    input_ids = torch.full((len(items), pmax), pad, dtype=torch.long, device=dev)
+    amask = torch.zeros((len(items), pmax), dtype=torch.long, device=dev)
+    pos = torch.ones((len(items), pmax), dtype=torch.long, device=dev)
+    for item_idx, (_rows, ids, _vit) in enumerate(items):
+        length = lengths[item_idx]
+        left = pmax - length
+        input_ids[item_idx, left:] = ids.to(dev)
+        amask[item_idx, left:] = 1
+        pos[item_idx, left:] = torch.arange(length, dtype=torch.long, device=dev)
+
+    visual_features = torch.cat([vit for _rows, _ids, vit in items], dim=0)
+    assert int((input_ids == img_tok).sum().item()) == visual_features.shape[0], \
+        "shared-prefix image-token count != supplied visual_features rows"
+
+    if debug:
+        group_sizes = [len(rows) for rows, _ids, _vit in items]
+        print(
+            "LA Flash hybrid shared prompt prefill "
+            f"groups={len(items)} group_sizes={group_sizes} prefix_lens={lengths}",
+            flush=True,
+        )
+
+    forward_mask, fallback_rows = _forward_attention_mask(
+        model, input_ids, amask, 0, mtp_window=False)
+    if debug and fallback_rows:
+        print(
+            "LA Flash hybrid shared prefill safe SDPA fallback "
+            f"query_rows={fallback_rows}",
+            flush=True,
+        )
+    forward_kwargs = dict(
+        input_ids=input_ids,
+        visual_features=visual_features,
+        image_token_index=img_tok,
+        attention_mask=forward_mask,
+        position_ids=pos,
+        past_key_values=None,
+        use_cache=True,
+    )
+    if isinstance(forward_mask, dict):
+        out = language_model_forward(model, **forward_kwargs, return_logits=False)
+    else:
+        out = model.language_model.model(**forward_kwargs)
+
+    real_tokens = sum(lengths)
+    shared_rows = sum(len(rows) for rows, _ids, _vit in items)
+    saved_tokens = sum((len(rows) - 1) * length for (rows, _ids, _vit), length in zip(items, lengths))
+    _record_prefill_stats(
+        stats,
+        rows=len(items),
+        q_len=pmax,
+        real_tokens=real_tokens,
+        shared_groups=len(items),
+        shared_rows=shared_rows,
+        saved_tokens=saved_tokens,
+    )
+
+    for item_idx, (rows, _ids, _vit) in enumerate(items):
+        prefix_len = lengths[item_idx]
+        prefix_kv = _unpack_stock_after_forward(out.past_key_values, item_idx, 0, prefix_len, 0, pmax)
+        for row in rows:
+            kv_rows[row] = prefix_kv
+            cached_lens[row] = prefix_len
+
+    return kv_rows, cached_lens
+
+
+@torch.no_grad()
+def _prefill_prompt_kv_rows(model, prompt_ids, vit_list, img_tok, pad, dev, mode, debug=False, stats=None):
+    """Return per-row prompt KV caches and cached lengths.
+
+    ``mode='none'`` preserves the legacy stock-like first MTP forward where the
+    whole prompt and the 6-token MTP window are forwarded together.  The split
+    prefill modes keep prompt KV clean before the scheduler batches only short
+    suffix/window forwards, which avoids ragged prompt+window masking in the
+    first decode step.
+    """
+    bsz = len(prompt_ids)
+    lengths = [int(p.numel()) for p in prompt_ids]
+    if mode == "none":
+        return [None] * bsz, [0] * bsz
+
+    base = model.language_model.model
+    if debug:
+        print(f"LA Flash hybrid prompt prefill mode={mode} rows={bsz} lengths={lengths}", flush=True)
+
+    if mode == "shared":
+        return _prefill_shared_prefix_kv_rows(
+            model, prompt_ids, vit_list, img_tok, pad, dev, stats=stats, debug=debug)
+
+    if mode == "per_row":
+        kv_rows = []
+        for b, ids in enumerate(prompt_ids):
+            ids = ids.to(dev).unsqueeze(0)
+            pos = torch.arange(ids.shape[1], dtype=torch.long, device=dev).unsqueeze(0)
+            out = base(
+                input_ids=ids,
+                visual_features=vit_list[b],
+                image_token_index=img_tok,
+                attention_mask=None,
+                position_ids=pos,
+                past_key_values=None,
+                use_cache=True,
+            )
+            kv_rows.append(out.past_key_values)
+            _record_prefill_stats(stats, rows=1, q_len=ids.shape[1], real_tokens=ids.shape[1])
+        return kv_rows, lengths
+
+    pmax = max(lengths)
+    input_ids = torch.full((bsz, pmax), pad, dtype=torch.long, device=dev)
+    amask = torch.zeros((bsz, pmax), dtype=torch.long, device=dev)
+    pos = torch.ones((bsz, pmax), dtype=torch.long, device=dev)
+    for b, ids in enumerate(prompt_ids):
+        left = pmax - lengths[b]
+        input_ids[b, left:] = ids.to(dev)
+        amask[b, left:] = 1
+        pos[b, left:] = torch.arange(lengths[b], dtype=torch.long, device=dev)
+
+    visual_features = torch.cat(vit_list, dim=0)
+    assert int((input_ids == img_tok).sum().item()) == visual_features.shape[0], \
+        "image-token count != supplied visual_features rows"
+    forward_mask, fallback_rows = _forward_attention_mask(
+        model, input_ids, amask, 0, mtp_window=False)
+    if debug and fallback_rows:
+        print(
+            "LA Flash hybrid batch prefill safe SDPA fallback "
+            f"query_rows={fallback_rows}",
+            flush=True,
+        )
+    forward_kwargs = dict(
+        input_ids=input_ids,
+        visual_features=visual_features,
+        image_token_index=img_tok,
+        attention_mask=forward_mask,
+        position_ids=pos,
+        past_key_values=None,
+        use_cache=True,
+    )
+    if isinstance(forward_mask, dict):
+        out = language_model_forward(model, **forward_kwargs, return_logits=False)
+    else:
+        out = base(**forward_kwargs)
+    _record_prefill_stats(stats, rows=bsz, q_len=pmax, real_tokens=sum(lengths))
+    kv_rows = [
+        _unpack_stock_after_forward(out.past_key_values, b, 0, lengths[b], 0, pmax)
+        for b in range(bsz)
+    ]
+    return kv_rows, lengths
+
+
+@torch.no_grad()
+def generate_batch_hybrid(pairs, temperature=README_TEMPERATURE, top_p=README_TOP_P, top_k=None,
+                          repetition_penalty=README_REPETITION_PENALTY,
+                          max_new_tokens=README_MAX_NEW_TOKENS, temps=None,
+                          debug=None, scheduler=None, group_size=None,
+                          vision_features=None, _stats=None):
+    """Batched stock-style LocateAnything-3B hybrid generation.
+
+    This mirrors ``model.generate(..., generation_mode='hybrid')``: each row
+    owns a full ``generated`` token stream plus a KV cache truncated to real
+    generated tokens before sampling.  MTP forwards
+    ``generated[cached_len:] + duplicate-last + mask*5``; AR forwards
+    ``generated[cached_len:]``.
+    """
+    tok, _, model = load()
+    san, hpat = _helpers()
+    tids = model.token_ids
+    img_tok = model.config.image_token_index
+    mask_tok = tids["default_mask_token_id"]
+    im_end = tids["im_end_token_id"]
+    pad = tok.pad_token_id if tok.pad_token_id is not None else im_end
+    dev = DEV
+
+    if not pairs:
+        return []
+    if temps is not None and len(temps) != len(pairs):
+        raise ValueError("temps must have the same length as pairs")
+    if vision_features is not None and len(vision_features) != len(pairs):
+        raise ValueError("vision_features must have the same length as pairs")
+    debug = _debug_enabled(debug)
+    scheduler = _hybrid_scheduler(scheduler)
+    group_size = _hybrid_group_size(group_size)
+    requested_attn = getattr(model, "_la_flash_requested_attn", ATTN_MODE)
+    use_magi = requested_attn == "magi"
+    prefill_mode = _hybrid_prefill_mode()
+    hold_max_steps = max(0, _env_int("LA_FLASH_HYBRID_HOLD_MAX_STEPS", 5))
+    adaptive_hold_mtp_max = max(0, _env_int("LA_FLASH_HYBRID_ADAPTIVE_HOLD_MTP_MAX", 3))
+    top_level_stats = _stats is None
+    if top_level_stats:
+        _stats = _new_hybrid_stats(
+            len(pairs), scheduler, group_size, hold_max_steps, adaptive_hold_mtp_max)
+        if os.environ.get("LA_FLASH_PLAN_STATS", "0") == "1":
+            model._la_flash_sparse_plan_stats = None
+    if group_size and len(pairs) > group_size:
+        outs = []
+        if debug:
+            print(
+                f"LA Flash hybrid grouped scheduling: total_rows={len(pairs)} "
+                f"group_size={group_size} scheduler={scheduler} hold_max_steps={hold_max_steps} "
+                f"adaptive_hold_mtp_max={adaptive_hold_mtp_max}",
+                flush=True,
+            )
+        for start in range(0, len(pairs), group_size):
+            end = min(start + group_size, len(pairs))
+            chunk_temps = temps[start:end] if temps is not None else None
+            chunk_vision_features = (
+                vision_features[start:end] if vision_features is not None else None
+            )
+            if debug:
+                print(f"LA Flash hybrid group rows=[{start}:{end}]", flush=True)
+            outs.extend(generate_batch_hybrid(
+                pairs[start:end],
+                temperature=temperature,
+                top_p=top_p,
+                top_k=top_k,
+                repetition_penalty=repetition_penalty,
+                max_new_tokens=max_new_tokens,
+                temps=chunk_temps,
+                debug=debug,
+                scheduler=scheduler,
+                group_size=0,
+                vision_features=chunk_vision_features,
+                _stats=_stats,
+            ))
+        if top_level_stats:
+            _set_last_hybrid_stats(_stats)
+        return outs
+
+    use_cached_tokenize = (
+        vision_features is not None
+        and os.environ.get("LA_FLASH_CACHE_TOKENIZE", "1") != "0"
+    )
+    if use_cached_tokenize:
+        try:
+            prompt_ids = [
+                _tokenize_cached_image(q, int(v.shape[0]), im=im)
+                for (im, q), v in zip(pairs, vision_features)
+            ]
+        except Exception as exc:
+            if os.environ.get("LA_FLASH_CACHE_TOKENIZE_STRICT", "0") == "1":
+                raise
+            if debug:
+                print(f"LA Flash cached tokenize fallback: {exc}", flush=True)
+            prompt_ids = [_tokenize(im, q) for im, q in pairs]
+    else:
+        prompt_ids = [_tokenize(im, q) for im, q in pairs]
+    vit_list = (
+        list(vision_features)
+        if vision_features is not None
+        else _encode_images([im for im, _ in pairs])
+    )
+    lengths = [int(p.numel()) for p in prompt_ids]
+    bsz = len(pairs)
+    _record_group_stats(_stats, bsz)
+
+    _set_llm_mode(model, requested_attn)
+
+    modes = ["mtp"] * bsz
+    finished = [False] * bsz
+    gen_ids = [[] for _ in range(bsz)]
+    full_ids = [list(ids.detach().cpu().tolist()) for ids in prompt_ids]
+    kv_rows, cached_lens = _prefill_prompt_kv_rows(
+        model, prompt_ids, vit_list, img_tok, pad, dev, prefill_mode, debug=debug, stats=_stats)
+    total_limits = [lengths[b] + max_new_tokens for b in range(bsz)]
+
+    row_temps = [float(temperature or 0.0)] * bsz if temps is None else [float(t or 0.0) for t in temps]
+
+    def run_ar(ar_rows, step_idx):
+        row_groups = _split_rows_by_kv_budget(ar_rows, kv_rows)
+        _record_kv_bucket_stats(_stats, row_groups, kv_rows)
+        for row_group in row_groups:
+            _step_stock_ar_rows(
+                model, san, tids, prompt_ids, kv_rows, row_group,
+                cached_lens, full_ids, gen_ids, modes, finished, total_limits,
+                pad, img_tok, row_temps, temperature, top_p, top_k,
+                repetition_penalty, dev, tok, debug, step_idx, use_magi, _stats,
+            )
+
+    def run_mtp(mtp_rows, step_idx):
+        if any(cached_lens[r] == 0 for r in mtp_rows) and any(cached_lens[r] > 0 for r in mtp_rows):
+            first_rows = [r for r in mtp_rows if cached_lens[r] == 0]
+            cached_rows = [r for r in mtp_rows if cached_lens[r] > 0]
+            if first_rows:
+                run_mtp(first_rows, step_idx)
+            if cached_rows:
+                run_mtp(cached_rows, step_idx)
+            return
+        row_groups = _split_rows_by_kv_budget(mtp_rows, kv_rows)
+        _record_kv_bucket_stats(_stats, row_groups, kv_rows)
+        if len(row_groups) > 1:
+            for row_group in row_groups:
+                run_mtp(row_group, step_idx)
+            return
+        _step_stock_mtp_rows(
+            model, san, hpat, tids, prompt_ids, kv_rows, mtp_rows,
+            cached_lens, full_ids, gen_ids, modes, finished, total_limits,
+            vit_list, pad, mask_tok, img_tok, row_temps, top_p, top_k,
+            repetition_penalty, dev, tok, debug, step_idx, use_magi, _stats,
+        )
+
+    def live_rows(mode):
+        return [b for b in range(bsz) if not finished[b] and modes[b] == mode]
+
+    step = 0
+    hold_steps = 0
+    while not all(finished) and step <= max_new_tokens:
+        step += 1
+        if _stats is not None:
+            _stats["decode_loops"] += 1
+        if scheduler == "hold_ar" and hold_max_steps > 0:
+            ar_rows = live_rows("ar")
+            mtp_rows = live_rows("mtp")
+            if ar_rows and mtp_rows and _stats is not None:
+                _stats["mixed_mode_cycles"] += 1
+            if ar_rows and (hold_steps < hold_max_steps or not mtp_rows):
+                if mtp_rows and _stats is not None:
+                    _stats["hold_ar_steps"] += 1
+                    _stats["hold_ar_held_mtp_rows"] += len(mtp_rows)
+                run_ar(ar_rows, step)
+                hold_steps += 1
+                continue
+            if mtp_rows:
+                if ar_rows and _stats is not None:
+                    _stats["hold_ar_limit_mtp_forwards"] += 1
+                run_mtp(mtp_rows, step)
+                hold_steps = 0
+            continue
+
+        if scheduler in {"ar_first", "pipeline", "adaptive"}:
+            ar_rows_at_loop_start = live_rows("ar")
+            mtp_rows_at_loop_start = live_rows("mtp")
+            mixed = bool(ar_rows_at_loop_start and mtp_rows_at_loop_start)
+            if mixed and _stats is not None:
+                _stats["mixed_mode_cycles"] += 1
+
+            if scheduler == "adaptive" and mixed and hold_max_steps > 0:
+                should_hold = len(mtp_rows_at_loop_start) <= adaptive_hold_mtp_max
+                if should_hold and hold_steps < hold_max_steps:
+                    if _stats is not None:
+                        _stats["adaptive_hold_cycles"] += 1
+                        _stats["hold_ar_steps"] += 1
+                        _stats["hold_ar_held_mtp_rows"] += len(mtp_rows_at_loop_start)
+                    run_ar(ar_rows_at_loop_start, step)
+                    hold_steps += 1
+                    continue
+
+            if ar_rows_at_loop_start:
+                if mixed and _stats is not None:
+                    if scheduler == "adaptive":
+                        _stats["adaptive_ar_first_cycles"] += 1
+                    else:
+                        _stats["ar_first_cycles"] += 1
+                run_ar(ar_rows_at_loop_start, step)
+
+            mtp_rows = live_rows("mtp")
+            if mtp_rows:
+                run_mtp(mtp_rows, step)
+                hold_steps = 0
+
+            if scheduler == "pipeline" and mtp_rows:
+                old_ar = set(ar_rows_at_loop_start)
+                new_ar_rows = [b for b in live_rows("ar") if b not in old_ar]
+                if new_ar_rows:
+                    if _stats is not None:
+                        _stats["pipeline_ar_after_mtp_cycles"] += 1
+                    run_ar(new_ar_rows, step)
+            continue
+
+        mtp_rows = live_rows("mtp")
+        ar_rows_at_loop_start = live_rows("ar")
+        if mtp_rows and ar_rows_at_loop_start and _stats is not None:
+            _stats["mixed_mode_cycles"] += 1
+        if mtp_rows:
+            run_mtp(mtp_rows, step)
+
+        ar_rows = [b for b in range(bsz) if not finished[b] and modes[b] == "ar"]
+        if mtp_rows and ar_rows and _stats is not None:
+            _stats["eager_mtp_then_ar_cycles"] += 1
+        if ar_rows:
+            run_ar(ar_rows, step)
+
+    outs = [
+        tok.decode(torch.tensor(gen_ids[b], dtype=torch.long, device=dev),
+                   skip_special_tokens=False) if gen_ids[b] else ""
+        for b in range(bsz)
+    ]
+    if top_level_stats:
+        if os.environ.get("LA_FLASH_PLAN_STATS", "0") == "1":
+            _stats["sparse_plan_stats"] = copy.deepcopy(
+                getattr(model, "_la_flash_sparse_plan_stats", None) or {}
+            )
+        _set_last_hybrid_stats(_stats)
+    return outs
+
+
+@torch.no_grad()
+def _step_stock_mtp_rows(model, san, hpat, tids, prompt_ids, kv_rows, rows,
+                         cached_lens, full_ids, gen_ids, modes, finished, total_limits,
+                         vit_list, pad, mask_tok, img_tok, row_temps, top_p, top_k,
+                         repetition_penalty, dev, tok, debug, step_idx, use_magi, stats=None):
+    kv, kvalid, old_lens, kmax = _pack_stock_kv_rows(kv_rows, rows, dev)
+    uncached_lens = [len(full_ids[r]) - cached_lens[r] for r in rows]
+    umax = max(uncached_lens)
+    seq_len = umax + N_FUTURE
+    _record_forward_stats(stats, "mtp", rows, seq_len, uncached_lens)
+
+    suf_ids = torch.full((len(rows), seq_len), pad, dtype=torch.long, device=dev)
+    suf_pos = torch.ones((len(rows), seq_len), dtype=torch.long, device=dev)
+    q_valid = torch.zeros((len(rows), seq_len), dtype=torch.long, device=dev)
+
+    for i, r in enumerate(rows):
+        uncached = full_ids[r][cached_lens[r] :]
+        left = umax - len(uncached)
+        if uncached:
+            suf_ids[i, left : left + len(uncached)] = torch.tensor(uncached, dtype=torch.long, device=dev)
+            suf_pos[i, left : left + len(uncached)] = torch.arange(
+                cached_lens[r], len(full_ids[r]), dtype=torch.long, device=dev)
+            q_valid[i, left : left + len(uncached)] = 1
+
+        rep = full_ids[r][-1]
+        cur_len = len(full_ids[r])
+        suf_ids[i, umax] = rep
+        suf_pos[i, umax] = cur_len - 1
+        q_valid[i, umax] = 1
+        for j in range(1, N_FUTURE):
+            suf_ids[i, umax + j] = mask_tok
+            suf_pos[i, umax + j] = cur_len + (j - 1)
+            q_valid[i, umax + j] = 1
+
+    full_mask = torch.cat([kvalid, q_valid], dim=1)
+
+    if debug:
+        forward_mask, fallback_rows = _forward_attention_mask(
+            model, suf_ids, full_mask, kmax, mtp_window=True, range_plan=True)
+        _print_debug_forward(
+            f"MTP step={step_idx}",
+            model,
+            tok,
+            suf_ids,
+            full_mask,
+            suf_pos,
+            past_len=kmax,
+            mtp_window=True,
+            extra={
+                "global_rows": rows,
+                "old_kv_lens": old_lens,
+                "cached_lens": [cached_lens[r] for r in rows],
+                "full_lens": [len(full_ids[r]) for r in rows],
+                "uncached_lens": uncached_lens,
+                "forward_attention_mask": _mask_desc(forward_mask),
+                "safe_sdpa_fallback_query_rows": fallback_rows,
+            },
+            attention_impl="magi" if use_magi else ATTN_MODE,
+        )
+    else:
+        forward_mask, _ = _forward_attention_mask(
+            model, suf_ids, full_mask, kmax, mtp_window=True, range_plan=True)
+    first_rows = [r for r in rows if cached_lens[r] == 0]
+    visual_features = None
+    if first_rows:
+        if first_rows != rows:
+            raise RuntimeError("mixed first/non-first MTP rows are not supported")
+        visual_features = torch.cat([vit_list[r] for r in rows], dim=0)
+        assert int((suf_ids == img_tok).sum().item()) == visual_features.shape[0], \
+            "image-token count != supplied visual_features rows"
+    out = language_model_forward(
+        model, input_ids=suf_ids, attention_mask=forward_mask,
+        position_ids=suf_pos, past_key_values=kv, use_cache=True,
+        visual_features=visual_features,
+        image_token_index=img_tok if visual_features is not None else None,
+        logits_slice=slice(-N_FUTURE, None))
+
+    for i, r in enumerate(rows):
+        kv_rows[r] = _unpack_stock_after_forward(
+            out.past_key_values, i, old_lens[i], uncached_lens[i], kmax, umax)
+        cached_lens[r] = len(full_ids[r])
+
+    wlogits = out.logits[:, -N_FUTURE:, :]
+    local_prompts = [prompt_ids[r] for r in rows]
+    local_gen = [gen_ids[r] for r in rows]
+    gen_pad = _pad_generated(local_prompts, local_gen, img_tok, dev)
+    per_row_temp = torch.tensor([row_temps[r] for r in rows], dtype=torch.float32, device=dev)
+
+    if BATCH_SAN:
+        x0_all, boxes_all = sample_tokens_batched(
+            wlogits, gen_pad, tids, per_row_temp,
+            repetition_penalty=repetition_penalty, top_p=top_p, top_k=top_k,
+            keep_k_avg=4, generation_mode="hybrid")
+
+    for i, r in enumerate(rows):
+        if finished[r]:
+            continue
+        if BATCH_SAN:
+            x0b, boxb = x0_all[i], boxes_all[i]
+        else:
+            gk = _mk_generate_kwargs(row_temps[r], top_p, top_k, repetition_penalty)
+            _, _, x0, box_avg = san(wlogits[i : i + 1], gen_pad[i : i + 1], tids, keep_k=5, **gk)
+            x0b, boxb = x0[0], box_avg[0]
+        nt = x0b if bool((boxb == 0).all()) else boxb
+        op = hpat(nt, tids, "hybrid")
+
+        toks = [int(t) for t in op["tokens"]]
+        for t in toks:
+            gen_ids[r].append(t)
+            full_ids[r].append(t)
+
+        if op["type"] == "im_end":
+            finished[r] = True
+        elif op["type"] == "error_box":
+            modes[r] = "ar"
+        if len(full_ids[r]) >= total_limits[r]:
+            finished[r] = True
+
+
+@torch.no_grad()
+def _step_stock_ar_rows(model, san, tids, prompt_ids, kv_rows, rows,
+                        cached_lens, full_ids, gen_ids, modes, finished, total_limits,
+                        pad, img_tok, row_temps, temperature, top_p, top_k,
+                        repetition_penalty, dev, tok, debug, step_idx, use_magi, stats=None):
+    kv, kvalid, old_lens, kmax = _pack_stock_kv_rows(kv_rows, rows, dev)
+    uncached_lens = [len(full_ids[r]) - cached_lens[r] for r in rows]
+    if any(n <= 0 for n in uncached_lens):
+        raise RuntimeError(f"AR rows have no uncached tokens: {rows}")
+    umax = max(uncached_lens)
+    _record_forward_stats(stats, "ar", rows, umax, uncached_lens)
+
+    suf_ids = torch.full((len(rows), umax), pad, dtype=torch.long, device=dev)
+    suf_pos = torch.ones((len(rows), umax), dtype=torch.long, device=dev)
+    q_valid = torch.zeros((len(rows), umax), dtype=torch.long, device=dev)
+
+    for i, r in enumerate(rows):
+        uncached = full_ids[r][cached_lens[r] :]
+        left = umax - len(uncached)
+        suf_ids[i, left:] = torch.tensor(uncached, dtype=torch.long, device=dev)
+        suf_pos[i, left:] = torch.arange(cached_lens[r], len(full_ids[r]), dtype=torch.long, device=dev)
+        q_valid[i, left:] = 1
+
+    full_mask = torch.cat([kvalid, q_valid], dim=1)
+    if debug:
+        forward_mask, fallback_rows = _forward_attention_mask(
+            model, suf_ids, full_mask, kmax, mtp_window=False, range_plan=True)
+        _print_debug_forward(
+            f"AR step={step_idx}",
+            model,
+            tok,
+            suf_ids,
+            full_mask,
+            suf_pos,
+            past_len=kmax,
+            mtp_window=False,
+            extra={
+                "global_rows": rows,
+                "old_kv_lens": old_lens,
+                "cached_lens": [cached_lens[r] for r in rows],
+                "full_lens": [len(full_ids[r]) for r in rows],
+                "uncached_lens": uncached_lens,
+                "forward_attention_mask": _mask_desc(forward_mask),
+                "safe_sdpa_fallback_query_rows": fallback_rows,
+            },
+            attention_impl="magi" if use_magi else ATTN_MODE,
+        )
+    else:
+        forward_mask, _ = _forward_attention_mask(
+            model, suf_ids, full_mask, kmax, mtp_window=False, range_plan=True)
+
+    out = language_model_forward(
+        model, input_ids=suf_ids, attention_mask=forward_mask,
+        position_ids=suf_pos, past_key_values=kv, use_cache=True,
+        logits_slice=slice(-1, None))
+
+    for i, r in enumerate(rows):
+        kv_rows[r] = _unpack_stock_after_forward(
+            out.past_key_values, i, old_lens[i], uncached_lens[i], kmax, umax)
+        cached_lens[r] = len(full_ids[r])
+
+    if AR_BATCH_SAN:
+        local_prompts = [prompt_ids[r] for r in rows]
+        local_gen = [gen_ids[r] for r in rows]
+        gen_pad = _pad_generated(local_prompts, local_gen, img_tok, dev)
+        per_row_temp = torch.tensor([row_temps[r] for r in rows], dtype=torch.float32, device=dev)
+        x0_all = sample_next_tokens_batched(
+            out.logits[:, -1:, :],
+            gen_pad,
+            per_row_temp,
+            repetition_penalty=repetition_penalty,
+            top_p=top_p,
+            top_k=top_k,
+        )
+
+    for i, r in enumerate(rows):
+        if AR_BATCH_SAN:
+            token_val = int(x0_all[i, 0].item())
+        else:
+            logits = out.logits[i : i + 1, -1:, :]
+            gen_pad = _pad_generated([prompt_ids[r]], [gen_ids[r]], img_tok, dev)
+            gk = _mk_generate_kwargs(temperature, top_p, top_k, repetition_penalty, row_temp=row_temps[r])
+            _, _, x0, _ = san(logits, gen_pad, tids, **gk)
+            token_val = int(x0[0, 0].item())
+        out_type = _classify_ar_token(token_val, tids)
+
+        gen_ids[r].append(token_val)
+        full_ids[r].append(token_val)
+
+        if out_type == "im_end":
+            finished[r] = True
+        elif out_type == "box_end_ar":
+            modes[r] = "mtp"
+
+        if len(full_ids[r]) >= total_limits[r]:
+            finished[r] = True
+
+
+def generate_batch_grouped_hybrid(groups, temperature=README_TEMPERATURE, top_p=README_TOP_P,
+                                  top_k=None, repetition_penalty=README_REPETITION_PENALTY,
+                                  max_new_tokens=README_MAX_NEW_TOKENS, temps=None,
+                                  debug=None, scheduler=None, group_size=None,
+                                  vision_features=None):
+    """Hybrid grouped API shape.
+
+    This preserves grouped return shape, but intentionally uses the generic
+    hybrid decoder rather than the fast engine's shared-prefix optimization.
+    """
+    flat = []
+    flat_vision_features = [] if vision_features is not None else None
+    counts = []
+    for group_idx, (im, queries) in enumerate(groups):
+        counts.append(len(queries))
+        flat.extend((im, q) for q in queries)
+        if flat_vision_features is not None:
+            flat_vision_features.extend([vision_features[group_idx]] * len(queries))
+
+    outs = generate_batch_hybrid(
+        flat, temperature=temperature, top_p=top_p, top_k=top_k,
+        repetition_penalty=repetition_penalty, max_new_tokens=max_new_tokens,
+        temps=temps, debug=debug, scheduler=scheduler, group_size=group_size,
+        vision_features=flat_vision_features)
+    res, offset = [], 0
+    for n in counts:
+        res.append(outs[offset : offset + n])
+        offset += n
+    return res
+
+
+__all__ = ["generate_batch_hybrid", "generate_batch_grouped_hybrid", "get_last_hybrid_stats"]

batch_utils/hybrid_runtime.py

@@ -0,0 +1,1842 @@
+"""Internal runtime support for the LocateAnything-3B hybrid batch decoder.
+
+This file keeps only the model-loading, tokenization, image-encoding, stock
+processor, and sample-token helpers that ``engine_hybrid.py`` needs.
+
+Important env knobs:
+  LA_FLASH_MODEL          HF repo id / local path of the model (default nvidia/LocateAnything-3B)
+  HF_HUB_OFFLINE=1    read the local HF cache only (no network); unset -> download on first use
+  LA_FLASH_ATTN           sdpa, eager, magi, or la_flash; la_flash uses FlashAttention sparse ranges
+  LA_FLASH_STRICT_ATTN    1 -> fail if the requested backend is unavailable;
+                      default 0 falls back to sdpa
+  LA_FLASH_VISION_ATTN    auto, flash_attention_2, sdpa, or eager (default auto)
+  LA_FLASH_HYBRID_PREFILL shared, none, per_row, or batch prompt KV prefill (default shared)
+  MTP_BATCH_VISION    0 -> per-image vision encode (default 1: batched when flash is present)
+  LA_FLASH_VISION_ENCODE_BATCH_SIZE
+                      max images per MoonViT encode micro-batch (default 8; <=0 disables limit)
+  MTP_BATCH_SAN       0 -> per-row logits/sample pipeline (default 1: batched over [B,6,V])
+  AR_BATCH_SAN        0 -> per-row AR sample pipeline (default 1: batched over [B,1,V])
+"""
+import inspect
+import os, warnings, importlib, torch
+from types import SimpleNamespace
+import numpy as np
+from transformers import AutoModel, AutoTokenizer, AutoProcessor
+
+
+# By default let transformers fetch the model on first use; set HF_HUB_OFFLINE=1 yourself
+# to read the local HF cache only (e.g. air-gapped / already-downloaded runs).
+MODEL = os.environ.get("LA_FLASH_MODEL", "nvidia/LocateAnything-3B")
+
+
+LLM_ATTN_MODES = ("sdpa", "eager", "magi", "la_flash")
+VISION_ATTN_MODES = ("auto", "flash_attention_2", "sdpa", "eager")
+
+
+def _normalize_attn_mode(value):
+    mode = (value or "sdpa").strip().lower().replace("-", "_")
+    aliases = {
+        "": "sdpa",
+        "manual": "eager",
+        "torch": "eager",
+        "torch_eager": "eager",
+        "torch_sdpa": "sdpa",
+        "scaled_dot_product_attention": "sdpa",
+        "flash": "la_flash",
+        "la_flash": "la_flash",
+        "kernel": "la_flash",
+        "cuda": "la_flash",
+        "range": "la_flash",
+        "range_attention": "la_flash",
+        "flex_flash": "magi",
+        "flex_flash_attention": "magi",
+        "flex_flash_attn": "magi",
+    }
+    mode = aliases.get(mode, mode)
+    if mode not in LLM_ATTN_MODES:
+        raise ValueError(
+            f"LA_FLASH_ATTN must be one of {', '.join(LLM_ATTN_MODES)}; got {value!r}"
+        )
+    return mode
+
+
+def _normalize_vision_attn_mode(value):
+    mode = (value or "auto").strip().lower().replace("-", "_")
+    aliases = {
+        "": "auto",
+        "flash": "flash_attention_2",
+        "flash_attention2": "flash_attention_2",
+        "fa2": "flash_attention_2",
+        "manual": "eager",
+    }
+    mode = aliases.get(mode, mode)
+    if mode not in VISION_ATTN_MODES:
+        raise ValueError(
+            f"LA_FLASH_VISION_ATTN must be one of {', '.join(VISION_ATTN_MODES)}; got {value!r}"
+        )
+    return mode
+
+
+ATTN_MODE = _normalize_attn_mode(os.environ.get("LA_FLASH_ATTN", "sdpa"))
+REMOTE_ATTN_MODE = "sdpa" if ATTN_MODE in {"la_flash", "magi"} else ATTN_MODE
+VISION_ATTN_MODE = _normalize_vision_attn_mode(os.environ.get("LA_FLASH_VISION_ATTN", "auto"))
+MAX_DIM = 1024
+DEV, DT = "cuda", torch.bfloat16
+N_FUTURE = 6                                            # = config.block_size (MTP window)
+_PROMPT = "Locate all the instances that matches the following description: "
+
+
+def _env_flag(name, default=False):
+    val = os.environ.get(name)
+    if val is None:
+        return default
+    return val.strip().lower() not in {"0", "false", "no", "off"}
+
+
+def _env_int(name):
+    val = os.environ.get(name)
+    if val is None or val.strip() == "":
+        return None
+    return int(val)
+
+
+def _strict_attn():
+    return _env_flag("LA_FLASH_STRICT_ATTN", False)
+
+
+def _fallback_to_sdpa(model, requested, reason):
+    if requested == "sdpa":
+        raise RuntimeError(f"LA_FLASH_ATTN=sdpa failed: {reason}") from reason
+    message = f"LA_FLASH_ATTN={requested} is unavailable; falling back to sdpa. Reason: {reason}"
+    if _strict_attn():
+        raise RuntimeError(message) from reason
+    warnings.warn(message)
+    _set_llm_mode(model, "sdpa")
+    model._la_flash_requested_attn_original = requested
+    model._la_flash_attn_fallback_reason = str(reason)
+    return "sdpa"
+
+
+# Optional compile for the shared Qwen2 core. This is off by default because the
+# hybrid scheduler already varies query/cache shapes and first-call compile cost is high.
+MTP_COMPILE = os.environ.get("MTP_COMPILE", "0") == "1"
+
+# Batch the MoonViT vision encode across a micro-batch's images: pack N images into ONE
+# extract_feature. With flash present, MoonViT's varlen cu_seqlens path is block-diagonal per
+# image and equivalent to per-image encode.
+# Without flash, sdpa builds a dense [1,S,S] mask -> O(S^2) N^2 -> per-image fallback (auto, see
+# _vision_is_flash). Default ON; set MTP_BATCH_VISION=0 to force per-image.
+BATCH_VISION = os.environ.get("MTP_BATCH_VISION", "1") == "1"
+_vision_encode_batch_size = _env_int("LA_FLASH_VISION_ENCODE_BATCH_SIZE")
+VISION_ENCODE_BATCH_SIZE = 8 if _vision_encode_batch_size is None else max(0, _vision_encode_batch_size)
+
+# Batch the per-row box-decode (sample_tokens): run the row-independent logits pipeline
+# (rep-penalty / per-row temperature / top_p / top_k / softmax / sample) ONCE over the whole
+# [B,6,V] step instead of B times on [1,6,V]; only the variable-length box assembly stays per-row.
+# Greedy is BIT-IDENTICAL to the per-row san (argmax, no RNG). Default ON; MTP_BATCH_SAN=0 -> per-row.
+BATCH_SAN = os.environ.get("MTP_BATCH_SAN", "1") == "1"
+
+# Batch the AR repair sampler over [B,1,V]. This shares the exact filtering
+# helpers with MTP batching but skips box/ref decoding, so it only replaces the
+# repeated stock one-token sample calls.  Sampling itself stays row-ordered by
+# default to preserve the stock RNG consumption pattern for AR repair.
+AR_BATCH_SAN = os.environ.get("AR_BATCH_SAN", "1") == "1"
+
+_tok = _proc = _model = None
+
+def _magi_diag():
+    lines = []
+    try:
+        import magi_attention
+        lines.append(f"magi_attention: OK file={getattr(magi_attention, '__file__', None)}")
+        lines.append(f"magi_attention.__version__={getattr(magi_attention, '__version__', '<missing>')}")
+    except Exception as e:
+        lines.append(f"magi_attention: FAIL {type(e).__name__}: {e}")
+        return "\n".join(lines)
+    try:
+        from magi_attention.functional.flex_flash_attn import flex_flash_attn_func
+        lines.append(f"magi_attention.functional.flex_flash_attn: OK func={flex_flash_attn_func}")
+    except Exception as e:
+        lines.append(f"magi_attention.functional.flex_flash_attn: FAIL {type(e).__name__}: {e}")
+    return "\n".join(lines)
+
+def _remote_magi_diag(model=None):
+    lines = []
+    try:
+        if model is not None:
+            mod = importlib.import_module(type(model.language_model.model).__module__)
+        else:
+            # Best effort: if the dynamic module is not imported yet this may fail;
+            # the post-load diagnostic below will still work.
+            mod = importlib.import_module("transformers_modules.LocateAnything-3B.modeling_qwen2")
+        lines.append(f"remote_qwen2_module={getattr(mod, '__file__', None)}")
+        lines.append(f"remote_qwen2._MAGI_AVAILABLE={getattr(mod, '_MAGI_AVAILABLE', '<missing>')!r}")
+        lines.append(f"remote_qwen2.flex_flash_attn_func={getattr(mod, 'flex_flash_attn_func', '<missing>')}")
+    except Exception as e:
+        lines.append(f"remote_qwen2: diagnostic failed {type(e).__name__}: {e}")
+    return "\n".join(lines)
+
+def _attn_class_diag(model):
+    try:
+        llm = model.language_model.model
+        classes = [type(layer.self_attn).__name__ for layer in llm.layers[:4]]
+        return (
+            f"llm._attn_implementation={getattr(llm, '_attn_implementation', None)!r}\n"
+            f"config._attn_implementation={getattr(llm.config, '_attn_implementation', None)!r}\n"
+            f"first_attn_classes={classes}"
+        )
+    except Exception as e:
+        return f"attention class diagnostic failed {type(e).__name__}: {e}"
+
+
+def _set_vision_attention_mode(model):
+    """Match HF's MoonViT policy: prefer flash_attention_2, then sdpa, then eager."""
+    vm = getattr(model, "vision_model", None)
+    if vm is None:
+        return None
+    mod = importlib.import_module(type(vm).__module__)
+    funcs = getattr(mod, "VL_VISION_ATTENTION_FUNCTIONS", {})
+    has_flash = getattr(mod, "flash_attn_varlen_func", None) is not None
+    requested = VISION_ATTN_MODE
+
+    if requested == "auto":
+        candidates = ("flash_attention_2", "sdpa", "eager")
+    else:
+        candidates = (requested, "flash_attention_2", "sdpa", "eager")
+
+    chosen = None
+    for candidate in candidates:
+        if candidate == "flash_attention_2" and not has_flash:
+            continue
+        if candidate in funcs:
+            chosen = candidate
+            break
+    if chosen is None:
+        raise RuntimeError("MoonViT has no supported attention implementation.")
+
+    if requested == "flash_attention_2" and chosen != "flash_attention_2":
+        warnings.warn("LA_FLASH_VISION_ATTN=flash_attention_2 requested but flash-attn is unavailable; "
+                      f"using {chosen}.")
+    elif requested not in {"auto", chosen}:
+        warnings.warn(f"LA_FLASH_VISION_ATTN={requested} is unavailable; using {chosen}.")
+
+    if hasattr(model.config, "vision_config"):
+        model.config.vision_config._attn_implementation = chosen
+    try:
+        vm.config._attn_implementation = chosen
+    except Exception:
+        pass
+    try:
+        for block in vm.encoder.blocks:
+            block.attn_implementation = chosen
+    except Exception as exc:
+        raise RuntimeError("Failed to configure MoonViT attention implementation.") from exc
+    model._la_flash_vision_attn = chosen
+    return chosen
+
+
+def load():
+    """Lazy model load with HF remote-code semantics plus release backends.
+
+    The text decoder is pinned to one of sdpa/eager/magi/la_flash. MoonViT is
+    configured independently and follows the HF policy: flash_attention_2 when
+    flash-attn is importable, otherwise sdpa, otherwise eager.
+    """
+    global _tok, _proc, _model
+    if _model is None:
+        _tok = AutoTokenizer.from_pretrained(MODEL, trust_remote_code=True)
+        _proc = AutoProcessor.from_pretrained(MODEL, trust_remote_code=True)
+        attn_impl = REMOTE_ATTN_MODE
+        if ATTN_MODE == "magi" and os.environ.get("LA_FLASH_DEBUG", "0") != "0":
+            print("LA Flash magi pre-load diagnostic:", flush=True)
+            print(_magi_diag(), flush=True)
+        _model = AutoModel.from_pretrained(MODEL, torch_dtype=DT, trust_remote_code=True,
+                                           attn_implementation=attn_impl).to(DEV).eval()
+        _set_vision_attention_mode(_model)
+        actual_attn = getattr(_model.language_model.model, "_attn_implementation", None)
+        if ATTN_MODE == "magi" and os.environ.get("LA_FLASH_DEBUG", "0") != "0":
+            print("LA Flash magi post-load diagnostic:", flush=True)
+            print(_remote_magi_diag(_model), flush=True)
+            print(_attn_class_diag(_model), flush=True)
+        if ATTN_MODE == "magi":
+            try:
+                qwen2_mod = importlib.import_module(type(_model.language_model.model).__module__)
+                if not getattr(qwen2_mod, "_MAGI_AVAILABLE", False):
+                    raise RuntimeError(
+                        "remote module reports _MAGI_AVAILABLE=False.\n"
+                        f"{_remote_magi_diag(_model)}\n{_magi_diag()}"
+                    )
+                first_attn = type(_model.language_model.model.layers[0].self_attn).__name__
+                if actual_attn != "sdpa" or first_attn != "_BatchedMagiAttention":
+                    _set_llm_mode(_model, "magi")
+                    actual_attn = getattr(_model.language_model.model, "_attn_implementation", None)
+                    first_attn = type(_model.language_model.model.layers[0].self_attn).__name__
+                    if os.environ.get("LA_FLASH_DEBUG", "0") != "0":
+                        print("LA Flash magi post-swap diagnostic:", flush=True)
+                        print(_attn_class_diag(_model), flush=True)
+                if actual_attn != "sdpa" or first_attn != "_BatchedMagiAttention":
+                    raise RuntimeError(
+                        "batched magi attention did not activate. "
+                        f"actual_attn={actual_attn!r}; first_attn={first_attn!r}; "
+                        f"{_remote_magi_diag(_model)}; {_attn_class_diag(_model)}"
+                    )
+                _model._la_flash_requested_attn = "magi"
+            except Exception as exc:
+                _fallback_to_sdpa(_model, "magi", exc)
+        else:
+            try:
+                _set_llm_mode(_model, ATTN_MODE)   # decode-safe mask plumbing for sdpa/eager/la_flash
+            except Exception as exc:
+                _fallback_to_sdpa(_model, ATTN_MODE, exc)
+        if MTP_COMPILE:
+            _maybe_compile(_model)
+    return _tok, _proc, _model
+
+
+def _maybe_compile(model):
+    """Compile the shared Qwen2Model core (base.forward). It backs BOTH prefill (called directly)
+    and decode (language_model.forward -> self.model). lm_head + MoonViT left eager. dynamic=True
+    so the varying decode S/kvlen don't trigger a recompile storm. No-op + warning if triton is
+    missing (inductor needs it on GPU). First call pays the compile cost (~42s warm / ~187s cold)."""
+    try:
+        import triton  # noqa: F401
+    except Exception:
+        warnings.warn("MTP_COMPILE set but triton is unavailable; running without torch.compile.")
+        return
+    import torch._dynamo as _dyn
+    _dyn.config.cache_size_limit = max(_dyn.config.cache_size_limit, 64)
+    base = model.language_model.model
+    if not getattr(base, "_mtp_compiled", False):
+        base.forward = torch.compile(base.forward, dynamic=True)
+        base._mtp_compiled = True
+
+
+def build_batched_magi_attention_class(mod):
+    """Build a Qwen2 attention subclass backed by Magi's flex_flash_attn.
+
+    The official LocateAnything ``Qwen2MagiAttention`` asserts ``bsz == 1`` and
+    relies on ``Qwen2Model._attn_implementation == "magi"`` to build a single
+    sample range plan.  For release batch inference the hybrid scheduler passes
+    a batched Magi range plan directly to this layer; a 4D-mask conversion path
+    remains as a compatibility fallback.
+    """
+    flex_flash_attn_func = getattr(mod, "flex_flash_attn_func", None)
+    if flex_flash_attn_func is None:
+        try:
+            from magi_attention.functional.flex_flash_attn import flex_flash_attn_func
+        except Exception as exc:
+            raise RuntimeError(
+                "LA_FLASH_ATTN=magi requires "
+                "magi_attention.functional.flex_flash_attn.flex_flash_attn_func."
+            ) from exc
+
+    FULL, CAUSAL = 0, 1
+    causal_plan_cache = {}
+    try:
+        magi_params = set(inspect.signature(flex_flash_attn_func).parameters)
+    except (TypeError, ValueError):
+        magi_params = set()
+    supports_disable_fwd_atomic = "disable_fwd_atomic_reduction" in magi_params
+
+    def _disjoint_q_ranges(q_ranges):
+        seen = set()
+        for start, end in q_ranges:
+            key = (int(start), int(end))
+            if key in seen:
+                return False
+            seen.add(key)
+        return True
+
+    def _plan_disjoint_q_ranges(plan):
+        cached = plan.get("_la_flash_disjoint_q_ranges")
+        if cached is not None:
+            return bool(cached)
+        q_ranges = plan["q_ranges"].detach().to(device="cpu", dtype=torch.int32).tolist()
+        disjoint = _disjoint_q_ranges(q_ranges)
+        try:
+            plan["_la_flash_disjoint_q_ranges"] = disjoint
+        except Exception:
+            pass
+        return disjoint
+
+    def _tensor_plan(q_ranges, k_ranges, types, device):
+        return {
+            "q_ranges": torch.tensor(q_ranges, dtype=torch.int32, device=device).contiguous(),
+            "k_ranges": torch.tensor(k_ranges, dtype=torch.int32, device=device).contiguous(),
+            "attn_type_map": torch.tensor(types, dtype=torch.int32, device=device).contiguous(),
+            "_la_flash_disjoint_q_ranges": _disjoint_q_ranges(q_ranges),
+        }
+
+    def _offset_plan(plan, q_offset, k_offset):
+        return (
+            (plan["q_ranges"] + int(q_offset)).tolist(),
+            (plan["k_ranges"] + int(k_offset)).tolist(),
+            plan["attn_type_map"].tolist(),
+        )
+
+    def _causal_plan(bsz, q_len, kv_seq_len, device):
+        key = (int(bsz), int(q_len), int(kv_seq_len), device.type, device.index)
+        cached = causal_plan_cache.get(key)
+        if cached is not None:
+            return cached
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            q_base = b * int(q_len)
+            k_base = b * int(kv_seq_len)
+            q_ranges.append([q_base, q_base + int(q_len)])
+            k_ranges.append([k_base, k_base + int(kv_seq_len)])
+            types.append(CAUSAL)
+        plan = _tensor_plan(q_ranges, k_ranges, types, device)
+        plan.update(
+            {
+                "flash_cu_seqlens_q": torch.arange(
+                    0,
+                    (int(bsz) + 1) * int(q_len),
+                    int(q_len),
+                    dtype=torch.int32,
+                    device=device,
+                ),
+                "flash_cu_seqlens_k": torch.arange(
+                    0,
+                    (int(bsz) + 1) * int(kv_seq_len),
+                    int(kv_seq_len),
+                    dtype=torch.int32,
+                    device=device,
+                ),
+                "flash_causal": True,
+            }
+        )
+        causal_plan_cache[key] = plan
+        return plan
+
+    def _row_segments(row):
+        idx = np.flatnonzero(row)
+        if idx.size == 0:
+            return ((0, 1),)
+        split = np.flatnonzero(np.diff(idx) > 1) + 1
+        starts = np.concatenate((idx[:1], idx[split]))
+        ends = np.concatenate((idx[split - 1], idx[-1:])) + 1
+        return tuple((int(s), int(e)) for s, e in zip(starts, ends))
+
+    def _visible_from_4d_mask(attention_mask, kv_seq_len):
+        mask = attention_mask[:, :, :, :kv_seq_len]
+        if mask.dtype == torch.bool:
+            return mask[:, 0].detach().to(device="cpu", dtype=torch.bool).contiguous()
+        mask_cpu = mask[:, 0].detach().to(device="cpu").contiguous()
+        if getattr(attention_mask, "_la_flash_visible_mask", False):
+            return (mask_cpu > 0).to(dtype=torch.bool)
+
+        max_value = float(mask_cpu.max().item()) if mask_cpu.numel() else 0.0
+        min_value = float(mask_cpu.min().item()) if mask_cpu.numel() else 0.0
+        if max_value > 0.0 and min_value >= 0.0:
+            return (mask_cpu > 0).to(dtype=torch.bool)
+        return (mask_cpu >= 0).to(dtype=torch.bool)
+
+    def _plan_from_visible_mask(attention_mask, bsz, q_len, kv_seq_len, device):
+        cache_key = (int(bsz), int(q_len), int(kv_seq_len), device.type, device.index)
+        cached = getattr(attention_mask, "_la_flash_magi_plan", None)
+        if cached is not None and cached[0] == cache_key:
+            return cached[1]
+
+        visible = _visible_from_4d_mask(attention_mask, int(kv_seq_len)).numpy()
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            q_base = b * int(q_len)
+            k_base = b * int(kv_seq_len)
+            run_start = 0
+            run_segments = _row_segments(visible[b, 0])
+            for q in range(1, int(q_len)):
+                segments = _row_segments(visible[b, q])
+                if segments == run_segments:
+                    continue
+                for start, end in run_segments:
+                    q_ranges.append([q_base + run_start, q_base + q])
+                    k_ranges.append([k_base + start, k_base + end])
+                    types.append(FULL)
+                run_start = q
+                run_segments = segments
+            for start, end in run_segments:
+                q_ranges.append([q_base + run_start, q_base + int(q_len)])
+                k_ranges.append([k_base + start, k_base + end])
+                types.append(FULL)
+
+        plan = _tensor_plan(q_ranges, k_ranges, types, device)
+        try:
+            attention_mask._la_flash_magi_plan = (cache_key, plan)
+        except Exception:
+            pass
+        return plan
+
+    def _plan_from_magi_dict(attention_mask, bsz, q_len, kv_seq_len, device):
+        if int(bsz) == 1:
+            return attention_mask
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            qs, ks, ts = _offset_plan(
+                attention_mask,
+                q_offset=b * int(q_len),
+                k_offset=b * int(kv_seq_len),
+            )
+            q_ranges.extend(qs)
+            k_ranges.extend(ks)
+            types.extend(ts)
+        return _tensor_plan(q_ranges, k_ranges, types, device)
+
+    def _magi_plan(attention_mask, bsz, q_len, kv_seq_len, device):
+        if isinstance(attention_mask, dict):
+            if attention_mask.get("_la_flash_batched", False):
+                return attention_mask
+            return _plan_from_magi_dict(attention_mask, bsz, q_len, kv_seq_len, device)
+        if attention_mask is None:
+            return _causal_plan(bsz, q_len, kv_seq_len, device)
+        if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+            raise ValueError(
+                f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, "
+                f"but is {attention_mask.size()}"
+            )
+        return _plan_from_visible_mask(attention_mask, bsz, q_len, kv_seq_len, device)
+
+    class _BatchedMagiAttention(mod.Qwen2Attention):
+        """MagiAttention path with true batch inference via packed token ranges."""
+
+        def forward(
+            self,
+            hidden_states: torch.Tensor,
+            attention_mask=None,
+            position_ids=None,
+            past_key_value=None,
+            output_attentions=False,
+            use_cache=False,
+            **kwargs,
+        ):
+            if output_attentions:
+                raise NotImplementedError("MagiAttention does not support output_attentions=True")
+
+            bsz, q_len, _ = hidden_states.size()
+            query_states = self.q_proj(hidden_states)
+            key_states = self.k_proj(hidden_states)
+            value_states = self.v_proj(hidden_states)
+
+            query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+            key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+            value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+            kv_seq_len = key_states.shape[-2]
+            if past_key_value is not None:
+                if self.layer_idx is None:
+                    raise ValueError(
+                        f"The cache structure has changed since version v4.36. If you are using "
+                        f"{self.__class__.__name__} for auto-regressive decoding with k/v caching, "
+                        "please initialize the attention class with a layer index."
+                    )
+                kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+
+            cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
+            query_states, key_states = mod.apply_rotary_pos_emb(
+                query_states, key_states, cos, sin, position_ids)
+
+            if past_key_value is not None:
+                cache_kwargs = {"sin": sin, "cos": cos}
+                key_states, value_states = past_key_value.update(
+                    key_states, value_states, self.layer_idx, cache_kwargs)
+
+            kv_seq_len = key_states.shape[-2]
+            plan = _magi_plan(attention_mask, bsz, q_len, kv_seq_len, query_states.device)
+            magi_extra_kwargs = {}
+            if supports_disable_fwd_atomic:
+                magi_extra_kwargs["disable_fwd_atomic_reduction"] = (
+                    (not self.training) and _plan_disjoint_q_ranges(plan)
+                )
+
+            query_states = query_states.transpose(1, 2).reshape(
+                bsz * q_len, self.num_heads, self.head_dim).contiguous()
+            key_states = key_states.transpose(1, 2).reshape(
+                bsz * kv_seq_len, self.num_key_value_heads, self.head_dim).contiguous()
+            value_states = value_states.transpose(1, 2).reshape(
+                bsz * kv_seq_len, self.num_key_value_heads, self.head_dim).contiguous()
+
+            attn_output, _ = flex_flash_attn_func(
+                query_states,
+                key_states,
+                value_states,
+                q_ranges=plan["q_ranges"],
+                k_ranges=plan["k_ranges"],
+                attn_type_map=plan["attn_type_map"],
+                softmax_scale=getattr(self, "softmax_scale", self.head_dim ** -0.5),
+                softcap=0.0,
+                deterministic=False,
+                **magi_extra_kwargs,
+            )
+            attn_output = attn_output.view(bsz, q_len, self.hidden_size)
+            attn_output = self.o_proj(attn_output)
+            return attn_output, None, past_key_value
+
+    return _BatchedMagiAttention
+
+
+def build_la_flash_attention_class(mod):
+    """Build a Qwen2 attention subclass backed by LA Flash sparse ranges."""
+    try:
+        from kernel_utils import is_available, range_attention
+    except Exception as exc:
+        raise RuntimeError(
+            "LA_FLASH_ATTN=la_flash requires kernel_utils and FlashAttention."
+        ) from exc
+    if not is_available():
+        raise RuntimeError(
+            "LA_FLASH_ATTN=la_flash requires flash_attn.flash_attn_varlen_func."
+        )
+
+    FULL, CAUSAL = 0, 1
+    causal_plan_cache = {}
+
+    def _tensor_plan(q_ranges, k_ranges, types, device):
+        max_q_len = max((int(end) - int(start) for start, end in q_ranges), default=0)
+        max_k_len = max((int(end) - int(start) for start, end in k_ranges), default=0)
+        plan = {
+            "q_ranges": torch.tensor(q_ranges, dtype=torch.int32, device=device).contiguous(),
+            "k_ranges": torch.tensor(k_ranges, dtype=torch.int32, device=device).contiguous(),
+            "attn_type_map": torch.tensor(types, dtype=torch.int32, device=device).contiguous(),
+            "max_q_len": max_q_len,
+            "max_k_len": max_k_len,
+        }
+        plan.update(_la_flash_group_plan_tensors(q_ranges, types, device))
+        return plan
+
+    def _offset_plan(plan, q_offset, k_offset):
+        return (
+            (plan["q_ranges"] + int(q_offset)).tolist(),
+            (plan["k_ranges"] + int(k_offset)).tolist(),
+            plan["attn_type_map"].tolist(),
+        )
+
+    def _causal_plan(bsz, q_len, kv_seq_len, device):
+        key = (int(bsz), int(q_len), int(kv_seq_len), device.type, device.index)
+        cached = causal_plan_cache.get(key)
+        if cached is not None:
+            return cached
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            q_base = b * int(q_len)
+            k_base = b * int(kv_seq_len)
+            q_ranges.append([q_base, q_base + int(q_len)])
+            k_ranges.append([k_base, k_base + int(kv_seq_len)])
+            types.append(CAUSAL)
+        plan = _tensor_plan(q_ranges, k_ranges, types, device)
+        plan.update(
+            {
+                "flash_cu_seqlens_q": torch.arange(
+                    0,
+                    (int(bsz) + 1) * int(q_len),
+                    int(q_len),
+                    dtype=torch.int32,
+                    device=device,
+                ),
+                "flash_cu_seqlens_k": torch.arange(
+                    0,
+                    (int(bsz) + 1) * int(kv_seq_len),
+                    int(kv_seq_len),
+                    dtype=torch.int32,
+                    device=device,
+                ),
+                "flash_causal": True,
+            }
+        )
+        causal_plan_cache[key] = plan
+        return plan
+
+    def _row_segments(row):
+        idx = np.flatnonzero(row)
+        if idx.size == 0:
+            return ((0, 1),)
+        split = np.flatnonzero(np.diff(idx) > 1) + 1
+        starts = np.concatenate((idx[:1], idx[split]))
+        ends = np.concatenate((idx[split - 1], idx[-1:])) + 1
+        return tuple((int(s), int(e)) for s, e in zip(starts, ends))
+
+    def _visible_from_4d_mask(attention_mask, kv_seq_len):
+        mask = attention_mask[:, :, :, :kv_seq_len]
+        if mask.dtype == torch.bool:
+            return mask[:, 0].detach().to(device="cpu", dtype=torch.bool).contiguous()
+        mask_cpu = mask[:, 0].detach().to(device="cpu").contiguous()
+        if getattr(attention_mask, "_la_flash_visible_mask", False):
+            return (mask_cpu > 0).to(dtype=torch.bool)
+
+        max_value = float(mask_cpu.max().item()) if mask_cpu.numel() else 0.0
+        min_value = float(mask_cpu.min().item()) if mask_cpu.numel() else 0.0
+        if max_value > 0.0 and min_value >= 0.0:
+            return (mask_cpu > 0).to(dtype=torch.bool)
+        return (mask_cpu >= 0).to(dtype=torch.bool)
+
+    def _prefix_len(row):
+        idx = np.flatnonzero(row)
+        if idx.size == 0:
+            return None
+        end = int(idx[-1]) + 1
+        if not bool(row[:end].all()) or bool(row[end:].any()):
+            return None
+        return end
+
+    def _causal_plan_from_visible(visible, bsz, q_len, kv_seq_len, device):
+        q_ranges, k_ranges, types = [], [], []
+        packed_flash = True
+        for b in range(int(bsz)):
+            first_len = _prefix_len(visible[b, 0])
+            if first_len is None:
+                return None
+            valid_len = int(first_len) + int(q_len) - 1
+            if valid_len < int(q_len) or valid_len > int(kv_seq_len):
+                return None
+            for q in range(int(q_len)):
+                row_len = _prefix_len(visible[b, q])
+                expected = valid_len - int(q_len) + q + 1
+                if row_len != expected:
+                    return None
+            q_base = b * int(q_len)
+            k_base = b * int(kv_seq_len)
+            q_ranges.append([q_base, q_base + int(q_len)])
+            k_ranges.append([k_base, k_base + valid_len])
+            types.append(CAUSAL)
+            packed_flash = packed_flash and valid_len == int(kv_seq_len)
+
+        plan = _tensor_plan(q_ranges, k_ranges, types, device)
+        plan["_la_flash_disjoint_q_ranges"] = True
+        if packed_flash:
+            plan.update(
+                {
+                    "flash_cu_seqlens_q": torch.arange(
+                        0,
+                        (int(bsz) + 1) * int(q_len),
+                        int(q_len),
+                        dtype=torch.int32,
+                        device=device,
+                    ),
+                    "flash_cu_seqlens_k": torch.arange(
+                        0,
+                        (int(bsz) + 1) * int(kv_seq_len),
+                        int(kv_seq_len),
+                        dtype=torch.int32,
+                        device=device,
+                    ),
+                    "flash_causal": True,
+                }
+            )
+        return plan
+
+    def _plan_from_visible_mask(attention_mask, bsz, q_len, kv_seq_len, device):
+        cache_key = (int(bsz), int(q_len), int(kv_seq_len), device.type, device.index, "la_flash")
+        cached = getattr(attention_mask, "_la_flash_range_plan", None)
+        if cached is not None and cached[0] == cache_key:
+            return cached[1]
+
+        visible = _visible_from_4d_mask(attention_mask, int(kv_seq_len)).numpy()
+        plan = _causal_plan_from_visible(visible, bsz, q_len, kv_seq_len, device)
+        if plan is not None:
+            try:
+                attention_mask._la_flash_range_plan = (cache_key, plan)
+            except Exception:
+                pass
+            return plan
+
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            q_base = b * int(q_len)
+            k_base = b * int(kv_seq_len)
+            run_start = 0
+            run_segments = _row_segments(visible[b, 0])
+            for q in range(1, int(q_len)):
+                segments = _row_segments(visible[b, q])
+                if segments == run_segments:
+                    continue
+                for start, end in run_segments:
+                    q_ranges.append([q_base + run_start, q_base + q])
+                    k_ranges.append([k_base + start, k_base + end])
+                    types.append(FULL)
+                run_start = q
+                run_segments = segments
+            for start, end in run_segments:
+                q_ranges.append([q_base + run_start, q_base + int(q_len)])
+                k_ranges.append([k_base + start, k_base + end])
+                types.append(FULL)
+
+        plan = _tensor_plan(q_ranges, k_ranges, types, device)
+        try:
+            attention_mask._la_flash_range_plan = (cache_key, plan)
+        except Exception:
+            pass
+        return plan
+
+    def _plan_from_magi_dict(attention_mask, bsz, q_len, kv_seq_len, device):
+        if int(bsz) == 1:
+            return attention_mask
+        q_ranges, k_ranges, types = [], [], []
+        for b in range(int(bsz)):
+            qs, ks, ts = _offset_plan(
+                attention_mask,
+                q_offset=b * int(q_len),
+                k_offset=b * int(kv_seq_len),
+            )
+            q_ranges.extend(qs)
+            k_ranges.extend(ks)
+            types.extend(ts)
+        return _tensor_plan(q_ranges, k_ranges, types, device)
+
+    def _range_plan(attention_mask, bsz, q_len, kv_seq_len, device):
+        if isinstance(attention_mask, dict):
+            if attention_mask.get("_la_flash_batched", False):
+                return attention_mask
+            return _plan_from_magi_dict(attention_mask, bsz, q_len, kv_seq_len, device)
+        if attention_mask is None:
+            return _causal_plan(bsz, q_len, kv_seq_len, device)
+        if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+            raise ValueError(
+                f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, "
+                f"but is {attention_mask.size()}"
+            )
+        return _plan_from_visible_mask(attention_mask, bsz, q_len, kv_seq_len, device)
+
+    class _LaFlashAttention(mod.Qwen2Attention):
+        """Range-plan attention path backed by FlashAttention sparse ranges."""
+
+        def forward(
+            self,
+            hidden_states: torch.Tensor,
+            attention_mask=None,
+            position_ids=None,
+            past_key_value=None,
+            output_attentions=False,
+            use_cache=False,
+            **kwargs,
+        ):
+            if output_attentions:
+                raise NotImplementedError("LA Flash attention does not support output_attentions=True")
+
+            bsz, q_len, _ = hidden_states.size()
+            query_states = self.q_proj(hidden_states)
+            key_states = self.k_proj(hidden_states)
+            value_states = self.v_proj(hidden_states)
+
+            query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+            key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+            value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+            kv_seq_len = key_states.shape[-2]
+            if past_key_value is not None:
+                if self.layer_idx is None:
+                    raise ValueError(
+                        f"The cache structure has changed since version v4.36. If you are using "
+                        f"{self.__class__.__name__} for auto-regressive decoding with k/v caching, "
+                        "please initialize the attention class with a layer index."
+                    )
+                kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+
+            cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
+            query_states, key_states = mod.apply_rotary_pos_emb(
+                query_states, key_states, cos, sin, position_ids)
+
+            if past_key_value is not None:
+                cache_kwargs = {"sin": sin, "cos": cos}
+                key_states, value_states = past_key_value.update(
+                    key_states, value_states, self.layer_idx, cache_kwargs)
+
+            kv_seq_len = key_states.shape[-2]
+            dense_backend = os.environ.get("LA_FLASH_DENSE_BACKEND", "sdpa").strip().lower()
+            if dense_backend == "sdpa" and not isinstance(attention_mask, dict):
+                dense_key_states = mod.repeat_kv(key_states, self.num_key_value_groups)
+                dense_value_states = mod.repeat_kv(value_states, self.num_key_value_groups)
+                if attention_mask is not None:
+                    if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+                        raise ValueError(
+                            f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, "
+                            f"but is {attention_mask.size()}"
+                        )
+                    query_for_sdpa = query_states.contiguous()
+                    key_for_sdpa = dense_key_states.contiguous()
+                    value_for_sdpa = dense_value_states.contiguous()
+                    is_causal = False
+                elif past_key_value is None:
+                    query_for_sdpa = query_states
+                    key_for_sdpa = dense_key_states
+                    value_for_sdpa = dense_value_states
+                    is_causal = bool(self.is_causal and q_len > 1)
+                else:
+                    query_for_sdpa = key_for_sdpa = value_for_sdpa = None
+                    is_causal = False
+                if query_for_sdpa is not None:
+                    attn_output = torch.nn.functional.scaled_dot_product_attention(
+                        query_for_sdpa,
+                        key_for_sdpa,
+                        value_for_sdpa,
+                        attn_mask=attention_mask,
+                        dropout_p=self.attention_dropout if self.training else 0.0,
+                        is_causal=is_causal,
+                    )
+                    attn_output = attn_output.transpose(1, 2).contiguous()
+                    attn_output = attn_output.reshape(bsz, q_len, self.hidden_size)
+                    attn_output = self.o_proj(attn_output)
+                    return attn_output, None, past_key_value
+
+            plan = _range_plan(attention_mask, bsz, q_len, kv_seq_len, query_states.device)
+
+            query_states = query_states.transpose(1, 2).reshape(
+                bsz * q_len, self.num_heads, self.head_dim).contiguous()
+            key_states = key_states.transpose(1, 2).reshape(
+                bsz * kv_seq_len, self.num_key_value_heads, self.head_dim).contiguous()
+            value_states = value_states.transpose(1, 2).reshape(
+                bsz * kv_seq_len, self.num_key_value_heads, self.head_dim).contiguous()
+
+            attn_output = range_attention(
+                query_states,
+                key_states,
+                value_states,
+                plan["q_ranges"],
+                plan["k_ranges"],
+                plan["attn_type_map"],
+                getattr(self, "softmax_scale", self.head_dim ** -0.5),
+                segment_offsets=plan.get("segment_offsets"),
+                group_q_ranges=plan.get("group_q_ranges"),
+                group_attn_type_map=plan.get("group_attn_type_map"),
+                max_q_len=plan.get("max_q_len"),
+                max_k_len=plan.get("max_k_len"),
+                flash_cu_seqlens_q=plan.get("flash_cu_seqlens_q"),
+                flash_cu_seqlens_k=plan.get("flash_cu_seqlens_k"),
+                flash_causal=plan.get("flash_causal"),
+                disjoint_q_ranges=plan.get("_la_flash_disjoint_q_ranges"),
+            )
+            attn_output = attn_output.view(bsz, q_len, self.hidden_size)
+            attn_output = self.o_proj(attn_output)
+            return attn_output, None, past_key_value
+
+    return _LaFlashAttention
+
+
+def _is_magi_plan(obj):
+    return isinstance(obj, dict) and {
+        "q_ranges",
+        "k_ranges",
+        "attn_type_map",
+    }.issubset(obj.keys())
+
+
+def _la_flash_group_plan_tensors(q_ranges, types, device):
+    """Group consecutive Magi range entries that share the same query span.
+
+    Magi-style plans may represent one query span with multiple disjoint key
+    spans.  LA Flash consumes those as one FlashAttention-backed softmax group.
+    """
+    if not q_ranges:
+        return {
+            "group_q_ranges": torch.empty((0, 2), dtype=torch.int32, device=device),
+            "segment_offsets": torch.zeros((1,), dtype=torch.int32, device=device),
+            "group_attn_type_map": torch.empty((0,), dtype=torch.int32, device=device),
+        }
+
+    grouped_q, grouped_types, offsets = [], [], [0]
+    last_q = None
+    last_type = None
+    for idx, (q_range, attn_type) in enumerate(zip(q_ranges, types)):
+        key = (int(q_range[0]), int(q_range[1]))
+        attn_type = int(attn_type)
+        if last_q is None:
+            grouped_q.append([key[0], key[1]])
+            grouped_types.append(attn_type)
+            last_q = key
+            last_type = attn_type
+            continue
+        if key == last_q and attn_type == last_type:
+            continue
+        offsets.append(idx)
+        grouped_q.append([key[0], key[1]])
+        grouped_types.append(attn_type)
+        last_q = key
+        last_type = attn_type
+    offsets.append(len(q_ranges))
+
+    return {
+        "group_q_ranges": torch.tensor(grouped_q, dtype=torch.int32, device=device).contiguous(),
+        "segment_offsets": torch.tensor(offsets, dtype=torch.int32, device=device).contiguous(),
+        "group_attn_type_map": torch.tensor(grouped_types, dtype=torch.int32, device=device).contiguous(),
+        "max_q_len": max((end - start for start, end in grouped_q), default=0),
+    }
+
+
+def _record_sparse_plan_stats(model, q_ranges, k_ranges, types):
+    if os.environ.get("LA_FLASH_PLAN_STATS", "0") != "1":
+        return
+    stats = getattr(model, "_la_flash_sparse_plan_stats", None)
+    if stats is None:
+        stats = {
+            "calls": 0,
+            "ranges": 0,
+            "q_tokens": 0,
+            "k_tokens": 0,
+            "max_q_len": 0,
+            "max_k_len": 0,
+            "full_ranges": 0,
+            "causal_ranges": 0,
+            "other_ranges": 0,
+        }
+        model._la_flash_sparse_plan_stats = stats
+    stats["calls"] += 1
+    stats["ranges"] += len(q_ranges)
+    for (q_start, q_end), (k_start, k_end), attn_type in zip(q_ranges, k_ranges, types):
+        q_len = int(q_end) - int(q_start)
+        k_len = int(k_end) - int(k_start)
+        stats["q_tokens"] += q_len
+        stats["k_tokens"] += k_len
+        stats["max_q_len"] = max(stats["max_q_len"], q_len)
+        stats["max_k_len"] = max(stats["max_k_len"], k_len)
+        attn_type = int(attn_type)
+        if attn_type == 0:
+            stats["full_ranges"] += 1
+        elif attn_type == 1:
+            stats["causal_ranges"] += 1
+        else:
+            stats["other_ranges"] += 1
+
+
+def build_magi_scheduler_ranges(model, attention_mask_2d, input_ids, past_len, mtp_window=False):
+    """Build batched Magi ranges directly from the hybrid scheduler mask.
+
+    The official Qwen2 SDPA dispatcher may optimize an all-valid 2D mask to
+    ``None`` before decoder layers see it. That is correct for plain causal
+    attention but loses LocateAnything's MTP generation-window rule. Building
+    ranges here keeps Magi batch inference exact and avoids per-layer dense
+    mask conversion.
+    """
+    requested_attn = getattr(model, "_la_flash_requested_attn", ATTN_MODE)
+    if requested_attn not in {"magi", "la_flash"}:
+        return None
+    if attention_mask_2d is None or not hasattr(attention_mask_2d, "dim") or attention_mask_2d.dim() != 2:
+        return None
+
+    bsz, q_len = int(input_ids.shape[0]), int(input_ids.shape[1])
+    key_len = int(attention_mask_2d.shape[1])
+    dev = input_ids.device
+    llm = model.language_model.model
+    block = int(getattr(llm, "block_size", N_FUTURE))
+    causal_attn = bool(getattr(llm, "causal_attn", False))
+    use_mtp_window = bool(mtp_window and q_len >= block and key_len >= block)
+    q0 = max(0, q_len - block)
+    k0 = max(0, key_len - block)
+    blocked_k = k0 - 1
+    past_len = int(past_len)
+
+    key_valid = attention_mask_2d.detach().to(device="cpu", dtype=torch.bool).contiguous().numpy()
+    key_idx = np.arange(key_len)
+    q_ranges, k_ranges, types = [], [], []
+    if not use_mtp_window:
+        causal_q_ranges, causal_k_ranges, causal_types = [], [], []
+        causal_fast_path = True
+        packed_flash = True
+        for b in range(bsz):
+            valid = np.flatnonzero(key_valid[b])
+            if valid.size == 0:
+                causal_fast_path = False
+                break
+            valid_len = int(valid[-1]) + 1
+            if valid_len < q_len or not bool(key_valid[b, :valid_len].all()) or bool(key_valid[b, valid_len:].any()):
+                causal_fast_path = False
+                break
+            packed_flash = packed_flash and valid_len == key_len
+            q_base = b * q_len
+            k_base = b * key_len
+            causal_q_ranges.append([q_base, q_base + q_len])
+            causal_k_ranges.append([k_base, k_base + valid_len])
+            causal_types.append(1)
+        if causal_fast_path:
+            plan = {
+                "q_ranges": torch.tensor(causal_q_ranges, dtype=torch.int32, device=dev).contiguous(),
+                "k_ranges": torch.tensor(causal_k_ranges, dtype=torch.int32, device=dev).contiguous(),
+                "attn_type_map": torch.tensor(causal_types, dtype=torch.int32, device=dev).contiguous(),
+                "max_q_len": q_len,
+                "max_k_len": max((end - start for start, end in causal_k_ranges), default=0),
+                "_la_flash_batched": True,
+                "_la_flash_disjoint_q_ranges": True,
+            }
+            if packed_flash:
+                plan.update(
+                    {
+                        "flash_cu_seqlens_q": torch.arange(
+                            0,
+                            (bsz + 1) * q_len,
+                            q_len,
+                            dtype=torch.int32,
+                            device=dev,
+                        ),
+                        "flash_cu_seqlens_k": torch.arange(
+                            0,
+                            (bsz + 1) * key_len,
+                            key_len,
+                            dtype=torch.int32,
+                            device=dev,
+                        ),
+                        "flash_causal": True,
+                    }
+                )
+            plan.update(_la_flash_group_plan_tensors(causal_q_ranges, causal_types, dev))
+            _record_sparse_plan_stats(model, causal_q_ranges, causal_k_ranges, causal_types)
+            return plan
+
+    def row_segments(row):
+        idx = np.flatnonzero(row)
+        if idx.size == 0:
+            return ((0, 1),)
+        split = np.flatnonzero(np.diff(idx) > 1) + 1
+        starts = np.concatenate((idx[:1], idx[split]))
+        ends = np.concatenate((idx[split - 1], idx[-1:])) + 1
+        return tuple((int(s), int(e)) for s, e in zip(starts, ends))
+
+    for b in range(bsz):
+        q_base = b * q_len
+        k_base = b * key_len
+        run_start = 0
+        run_segments = None
+        if use_mtp_window and not causal_attn:
+            prefix_q_len = q0
+            prefix_k_end = past_len + prefix_q_len
+            prefix_ok = (
+                prefix_q_len > 0
+                and prefix_k_end <= key_len
+                and bool(key_valid[b, :prefix_k_end].all())
+            )
+            window_prefix_ok = blocked_k <= 0 or bool(key_valid[b, :blocked_k].all())
+            window_ok = bool(key_valid[b, k0:key_len].all())
+            if prefix_ok:
+                q_ranges.append([q_base, q_base + prefix_q_len])
+                k_ranges.append([k_base, k_base + prefix_k_end])
+                types.append(1)
+                run_start = prefix_q_len
+            if run_start == prefix_q_len and prefix_q_len < q_len and window_prefix_ok and window_ok:
+                if blocked_k > 0:
+                    q_ranges.append([q_base + prefix_q_len, q_base + q_len])
+                    k_ranges.append([k_base, k_base + blocked_k])
+                    types.append(0)
+                q_ranges.append([q_base + prefix_q_len, q_base + q_len])
+                k_ranges.append([k_base + k0, k_base + key_len])
+                types.append(0)
+                continue
+
+        for q in range(run_start, q_len):
+            visible = key_valid[b] & (key_idx <= q + past_len)
+            if use_mtp_window and q >= q0:
+                if not causal_attn:
+                    visible = visible.copy()
+                    visible[k0:key_len] = key_valid[b, k0:key_len]
+                if blocked_k >= 0:
+                    if visible.base is None:
+                        visible[blocked_k] = False
+                    else:
+                        visible = visible.copy()
+                        visible[blocked_k] = False
+            segments = row_segments(visible)
+            if run_segments is None:
+                run_segments = segments
+                continue
+            if segments == run_segments:
+                continue
+            for start, end in run_segments:
+                q_ranges.append([q_base + run_start, q_base + q])
+                k_ranges.append([k_base + start, k_base + end])
+                types.append(0)
+            run_start = q
+            run_segments = segments
+        for start, end in run_segments:
+            q_ranges.append([q_base + run_start, q_base + q_len])
+            k_ranges.append([k_base + start, k_base + end])
+            types.append(0)
+
+    seen_q_ranges = set()
+    disjoint_q_ranges = True
+    for start, end in q_ranges:
+        key = (int(start), int(end))
+        if key in seen_q_ranges:
+            disjoint_q_ranges = False
+            break
+        seen_q_ranges.add(key)
+
+    plan = {
+        "q_ranges": torch.tensor(q_ranges, dtype=torch.int32, device=dev).contiguous(),
+        "k_ranges": torch.tensor(k_ranges, dtype=torch.int32, device=dev).contiguous(),
+        "attn_type_map": torch.tensor(types, dtype=torch.int32, device=dev).contiguous(),
+        "max_q_len": max((end - start for start, end in q_ranges), default=0),
+        "max_k_len": max((end - start for start, end in k_ranges), default=0),
+        "_la_flash_batched": True,
+        "_la_flash_disjoint_q_ranges": disjoint_q_ranges,
+    }
+    plan.update(_la_flash_group_plan_tensors(q_ranges, types, dev))
+    _record_sparse_plan_stats(model, q_ranges, k_ranges, types)
+    return plan
+
+
+def _direct_base_forward(
+    base,
+    input_ids=None,
+    visual_features=None,
+    image_token_index=None,
+    attention_mask=None,
+    position_ids=None,
+    past_key_values=None,
+    inputs_embeds=None,
+    use_cache=None,
+    output_attentions=None,
+    output_hidden_states=None,
+    return_dict=None,
+):
+    mod = importlib.import_module(type(base).__module__)
+    output_attentions = output_attentions if output_attentions is not None else base.config.output_attentions
+    output_hidden_states = (
+        output_hidden_states if output_hidden_states is not None else base.config.output_hidden_states
+    )
+    use_cache = use_cache if use_cache is not None else base.config.use_cache
+
+    if input_ids is not None and inputs_embeds is not None:
+        raise ValueError("You cannot specify both decoder_input_ids and decoder_inputs_embeds at the same time")
+    if input_ids is not None:
+        batch_size, seq_length = input_ids.shape
+    elif inputs_embeds is not None:
+        batch_size, seq_length, _ = inputs_embeds.shape
+    else:
+        raise ValueError("You have to specify either decoder_input_ids or decoder_inputs_embeds")
+
+    past_key_values_length = 0
+    use_legacy_cache = False
+    if use_cache:
+        Cache = getattr(mod, "Cache")
+        DynamicCache = getattr(mod, "DynamicCache")
+        use_legacy_cache = not isinstance(past_key_values, Cache)
+        if use_legacy_cache:
+            if past_key_values is None:
+                past_key_values = DynamicCache()
+            else:
+                past_key_values = DynamicCache.from_legacy_cache(past_key_values)
+        past_key_values_length = past_key_values.get_seq_length()
+
+    if position_ids is None:
+        dev = input_ids.device if input_ids is not None else inputs_embeds.device
+        position_ids = torch.arange(
+            past_key_values_length,
+            seq_length + past_key_values_length,
+            dtype=torch.long,
+            device=dev,
+        ).unsqueeze(0).view(-1, seq_length)
+    else:
+        position_ids = position_ids.view(-1, seq_length).long()
+
+    if inputs_embeds is None:
+        inputs_embeds = base.image_processing(input_ids, visual_features, image_token_index)
+
+    hidden_states = inputs_embeds
+    all_hidden_states = () if output_hidden_states else None
+    all_self_attns = () if output_attentions else None
+    next_decoder_cache = None
+
+    for decoder_layer in base.layers:
+        if output_hidden_states:
+            all_hidden_states += (hidden_states,)
+        layer_outputs = decoder_layer(
+            hidden_states,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_value=past_key_values,
+            output_attentions=output_attentions,
+            use_cache=use_cache,
+        )
+        hidden_states = layer_outputs[0]
+        if use_cache:
+            next_decoder_cache = layer_outputs[2 if output_attentions else 1]
+        if output_attentions:
+            all_self_attns += (layer_outputs[1],)
+
+    hidden_states = base.norm(hidden_states)
+    if output_hidden_states:
+        all_hidden_states += (hidden_states,)
+    next_cache = None
+    if use_cache:
+        next_cache = next_decoder_cache.to_legacy_cache() if use_legacy_cache else next_decoder_cache
+    return SimpleNamespace(
+        last_hidden_state=hidden_states,
+        past_key_values=next_cache,
+        hidden_states=all_hidden_states,
+        attentions=all_self_attns,
+    )
+
+
+def language_model_forward(model, **kwargs):
+    """Forward through the text LM, bypassing official dense-mask prep for sparse plans."""
+    lm = model.language_model
+    return_logits = kwargs.pop("return_logits", True)
+    logits_slice = kwargs.pop("logits_slice", None)
+    attention_mask = kwargs.get("attention_mask")
+    use_direct_sparse = (
+        getattr(model, "_la_flash_requested_attn", ATTN_MODE) in {"magi", "la_flash"}
+        and _is_magi_plan(attention_mask)
+    )
+    if not use_direct_sparse:
+        return lm(**kwargs)
+
+    labels = kwargs.pop("labels", None)
+    if labels is not None:
+        raise NotImplementedError("labels are not supported in the direct sparse-plan decode forward")
+    output_attentions = kwargs.get("output_attentions", None)
+    output_hidden_states = kwargs.get("output_hidden_states", None)
+    base_out = _direct_base_forward(lm.model, **kwargs)
+    logits = None
+    if return_logits:
+        hidden_states = base_out.last_hidden_state
+        if logits_slice is not None:
+            hidden_states = hidden_states[:, logits_slice, :]
+        logits = lm.lm_head(hidden_states).float()
+    return SimpleNamespace(
+        logits=logits,
+        past_key_values=base_out.past_key_values,
+        hidden_states=base_out.hidden_states if output_hidden_states else None,
+        attentions=base_out.attentions if output_attentions else None,
+    )
+
+
+_EagerCls = _SdpaCls = _LaFlashCls = _MagiCls = None
+def _attn_classes(mode=None):
+    """Attention classes from the dynamic Qwen2 remote module.
+
+    The official Qwen2Model mask dispatcher only implements ``sdpa`` and
+    single-row ``magi``.  Eager, LA Flash, and batched Magi inference
+    therefore swap the layer class while keeping the model's mask dispatcher
+    pinned to ``sdpa``.
+    """
+    global _EagerCls, _SdpaCls, _LaFlashCls, _MagiCls
+    mode = _normalize_attn_mode(mode) if mode is not None else None
+    if _SdpaCls is None:
+        mod = importlib.import_module(type(_model.language_model.model).__module__)
+        _EagerCls = mod.Qwen2Attention
+        _SdpaCls = mod.Qwen2SdpaAttention
+    else:
+        mod = importlib.import_module(type(_model.language_model.model).__module__)
+    if (mode is None or mode == "la_flash") and _LaFlashCls is None:
+        _LaFlashCls = build_la_flash_attention_class(mod)
+    if (mode is None or mode == "magi") and _MagiCls is None:
+        _MagiCls = build_batched_magi_attention_class(mod) if getattr(mod, "_MAGI_AVAILABLE", False) else None
+    return _EagerCls, _SdpaCls, _LaFlashCls, _MagiCls
+
+def _set_llm_mode(model, mode):
+    """Swap every Qwen2 decoder layer's attention class.
+
+    Release backends keep ``Qwen2Model._attn_implementation='sdpa'`` so the
+    official Qwen2 mask dispatcher stays available for dense-mask modes. The
+    local ``la_flash`` and batched ``magi`` wrappers can also consume scheduler-built
+    sparse plans directly, avoiding repeated per-layer dense mask conversion.
+    """
+    mode = _normalize_attn_mode(mode)
+    eager, sdpa, la_flash, magi = _attn_classes(mode)
+    impl = "sdpa"
+    if mode == "sdpa":
+        cls = sdpa
+    elif mode == "eager":
+        cls = eager
+    elif mode == "la_flash":
+        cls = la_flash
+    elif mode == "magi":
+        if magi is None:
+            raise RuntimeError("MagiAttention is unavailable in the current Python environment.")
+        cls = magi
+    else:
+        raise ValueError(f"unknown LLM attention mode: {mode}")
+    llm = model.language_model.model
+    for lyr in llm.layers:
+        lyr.self_attn.__class__ = cls
+        if mode == "magi":
+            lyr.self_attn.softmax_scale = lyr.self_attn.head_dim ** -0.5
+    llm._attn_implementation = impl
+    llm.config._attn_implementation = llm._attn_implementation
+    if hasattr(model.config, "text_config"):
+        model.config.text_config._attn_implementation = llm._attn_implementation
+    model.config._attn_implementation = llm._attn_implementation
+    model._la_flash_requested_attn = mode
+
+_st = _hp = None
+def _helpers():
+    """The model's own sample_tokens / handle_pattern (the exact box decoders)."""
+    global _st, _hp
+    if _st is None:
+        m = importlib.import_module(type(load()[2]).__module__)
+        _st, _hp = m.sample_tokens, m.handle_pattern
+    return _st, _hp
+
+
+_gu = None
+def _gen_utils():
+    """The model's generate_utils module (apply_repetition_penalty / top_p_logits / top_k_logits /
+    decode_bbox_avg / decode_ref / dists) -- the pieces sample_tokens_batched reuses verbatim."""
+    global _gu
+    if _gu is None:
+        m = importlib.import_module(type(load()[2]).__module__)
+        _gu = importlib.import_module(m.sample_tokens.__module__)
+    return _gu
+
+
+def _env_float(name, default):
+    val = os.environ.get(name)
+    if val is None or val.strip() == "":
+        return float(default)
+    return float(val)
+
+
+def _coord_fallback_mode():
+    mode = os.environ.get("LA_FLASH_COORD_FALLBACK_MODE", "legacy").strip().lower().replace("-", "_")
+    aliases = {
+        "": "legacy",
+        "official": "legacy",
+        "range": "legacy",
+        "spread": "legacy",
+        "none": "off",
+        "disable": "off",
+        "disabled": "off",
+        "entropy_variance": "uncertainty",
+        "entropy_var": "uncertainty",
+        "ent_var": "uncertainty",
+        "entropy_std": "uncertainty",
+    }
+    mode = aliases.get(mode, mode)
+    if mode not in {"legacy", "uncertainty", "off"}:
+        raise ValueError(
+            "LA_FLASH_COORD_FALLBACK_MODE must be one of legacy, uncertainty, off"
+        )
+    return mode
+
+
+def _coord_uncertainty_threshold(coord_start_token_id, coord_end_token_id):
+    """Return the coord uncertainty threshold in raw coord-token units.
+
+    Backward-compatible behavior:
+    - LA_FLASH_COORD_UNCERTAINTY_THRESH > 1 is treated as raw coord-token RMSE.
+    - LA_FLASH_COORD_UNCERTAINTY_THRESH <= 1 is treated as normalized by coord span.
+    - LA_FLASH_COORD_UNCERTAINTY_NORM_THRESH is an explicit normalized override.
+    """
+    coord_span = max(float(coord_end_token_id - coord_start_token_id + 1), 1.0)
+    norm_val = os.environ.get("LA_FLASH_COORD_UNCERTAINTY_NORM_THRESH")
+    if norm_val is not None and norm_val.strip() != "":
+        return float(norm_val) * coord_span
+
+    val = os.environ.get("LA_FLASH_COORD_UNCERTAINTY_THRESH")
+    if val is None or val.strip() == "":
+        return 20.0
+    threshold = float(val)
+    if 0.0 < threshold <= 1.0:
+        return threshold * coord_span
+    return threshold
+
+
+def _decode_bbox_with_uncertainty(logits, probs, token_ids, keep_k=4, generation_mode="hybrid"):
+    """Decode an MTP box with configurable coord uncertainty fallback.
+
+    The default mode is the official LocateAnything rule.  ``uncertainty`` keeps
+    the same frame checks and top-k coord selection, but uses one scalar
+    criterion per coordinate: the posterior RMSE of committing to the current
+    MAP coordinate among valid coord candidates.  This is the Bayes risk under
+    squared coordinate error, so probabilities and token distances are folded
+    into one threshold in coordinate-token units.
+    """
+    gu = _gen_utils()
+    mode = _coord_fallback_mode()
+    if mode == "legacy" or generation_mode != "hybrid":
+        return gu.decode_bbox_avg(logits, probs, token_ids, keep_k=keep_k, generation_mode=generation_mode)
+
+    coord_start_token_id = token_ids["coord_start_token_id"]
+    coord_end_token_id = token_ids["coord_end_token_id"]
+    box_start_token_id = token_ids["box_start_token_id"]
+    box_end_token_id = token_ids["box_end_token_id"]
+    none_token_id = token_ids["none_token_id"]
+    null_token_id = token_ids["null_token_id"]
+    device = logits.device
+
+    box_type = gu.is_valid_box_frame(
+        probs,
+        token_ids,
+        start_thresh=_env_float("LA_FLASH_COORD_BOX_START_THRESH", 0.7),
+        end_thresh=_env_float("LA_FLASH_COORD_BOX_END_THRESH", 0.2),
+        topk=keep_k,
+    )
+    if box_type == "empty_box":
+        return torch.tensor([
+            box_start_token_id,
+            none_token_id,
+            box_end_token_id,
+            null_token_id,
+            null_token_id,
+            null_token_id,
+        ], dtype=torch.long, device=device)
+    if box_type == "illegal_box":
+        return None
+
+    pos_probs, pos_ids = torch.topk(probs[1:5], k=keep_k, dim=-1)
+    valid = (pos_ids >= coord_start_token_id) & (pos_ids <= coord_end_token_id)
+    has_valid = valid.any(dim=-1)
+    if not has_valid.all():
+        return None
+
+    first_valid_idx = valid.long().argmax(dim=-1, keepdim=True)
+    first_valid_ids = pos_ids.gather(-1, first_valid_idx).squeeze(-1)
+    if mode == "off":
+        final_coords = first_valid_ids
+    else:
+        valid_counts = valid.sum(dim=-1)
+        valid_probs = torch.where(valid, pos_probs, torch.zeros_like(pos_probs))
+        valid_mass = valid_probs.sum(dim=-1).clamp_min(1e-12)
+        weights = valid_probs / valid_mass.unsqueeze(-1)
+        coord_values = (pos_ids - coord_start_token_id).to(dtype=torch.float32)
+        map_coord = (first_valid_ids - coord_start_token_id).to(dtype=torch.float32)
+        uncertainty = (weights * (coord_values - map_coord.unsqueeze(-1)).pow(2)).sum(dim=-1).sqrt()
+        is_abnormal = (
+            (valid_counts > 1)
+            & (uncertainty > _coord_uncertainty_threshold(coord_start_token_id, coord_end_token_id))
+        )
+        final_coords = torch.where(is_abnormal, torch.tensor(0, device=device), first_valid_ids)
+
+    start_t = torch.tensor([box_start_token_id], dtype=final_coords.dtype, device=device)
+    end_t = torch.tensor([box_end_token_id], dtype=final_coords.dtype, device=device)
+    return torch.cat([start_t, final_coords, end_t])
+
+
+def _apply_repetition_penalty_lowmem(logits, generated, repetition_penalty):
+    """Apply the stock repetition penalty without allocating a [B, S, V] mask."""
+    if repetition_penalty == 1.0:
+        return logits
+    _, _, vocab_size = logits.shape
+    for row in range(logits.shape[0]):
+        valid_tokens = generated[row].unique()
+        valid_tokens = valid_tokens[(valid_tokens >= 0) & (valid_tokens < vocab_size)]
+        if valid_tokens.numel() == 0:
+            continue
+        row_logits = logits[row, :, valid_tokens]
+        logits[row, :, valid_tokens] = torch.where(
+            row_logits > 0,
+            row_logits / repetition_penalty,
+            row_logits * repetition_penalty,
+        )
+    return logits
+
+
+def _finite_logit_bounds(dtype):
+    finfo = torch.finfo(dtype)
+    return finfo.min, finfo.max
+
+
+def _finite_logits(logits):
+    if not logits.dtype.is_floating_point:
+        logits = logits.float()
+    min_val, max_val = _finite_logit_bounds(logits.dtype)
+    return torch.nan_to_num(logits, nan=min_val, posinf=max_val, neginf=min_val)
+
+
+def _finite_logits_(logits):
+    if not logits.dtype.is_floating_point:
+        return logits.float()
+    min_val, max_val = _finite_logit_bounds(logits.dtype)
+    return logits.nan_to_num_(nan=min_val, posinf=max_val, neginf=min_val)
+
+
+def _top_p_logits_slice_(logits, top_p):
+    sorted_logits, sorted_indices = torch.sort(logits, descending=True)
+    cumulative_probs = torch.cumsum(torch.softmax(sorted_logits, dim=-1), dim=-1)
+    sorted_indices_to_remove = cumulative_probs > top_p
+    sorted_indices_to_remove[..., 1:] = sorted_indices_to_remove[..., :-1].clone()
+    sorted_indices_to_remove[..., 0] = False
+
+    remove = torch.zeros_like(logits, dtype=torch.bool, device=logits.device)
+    remove.scatter_(-1, sorted_indices, sorted_indices_to_remove)
+    logits.masked_fill_(remove, torch.finfo(logits.dtype).min)
+    return logits
+
+
+def _top_p_logits_(logits, top_p):
+    """In-place nucleus filtering with bounded sort workspace.
+
+    The MTP sampler uses logits shaped ``[B, 6, V]``.  Top-p is independent for
+    each row and each future position, so filtering one position at a time keeps
+    the expensive sorted-index workspace at ``[B, V]`` instead of ``[B, 6, V]``.
+    """
+    if logits.dim() == 3 and logits.shape[1] > 1:
+        for pos in range(logits.shape[1]):
+            _top_p_logits_slice_(logits[:, pos, :], top_p)
+        return logits
+    return _top_p_logits_slice_(logits, top_p)
+
+
+def _top_k_logits_(logits, top_k):
+    """In-place top-k filtering mirroring generate_utils.top_k_logits."""
+    top_k = min(int(top_k), logits.size(-1))
+    threshold = torch.topk(logits, top_k)[0][..., -1, None]
+    logits.masked_fill_(logits < threshold, torch.finfo(logits.dtype).min)
+    return logits
+
+
+def _safe_probs(filtered_logits):
+    """Softmax with CUDA-multinomial-safe cleanup and row-wise argmax fallback."""
+    filtered_logits = _finite_logits(filtered_logits)
+    probs = torch.softmax(filtered_logits, dim=-1, dtype=torch.float32)
+    probs = torch.nan_to_num(probs, nan=0.0, posinf=0.0, neginf=0.0).clamp_min_(0.0)
+    row_sum = probs.sum(dim=-1, keepdim=True)
+    bad = (~torch.isfinite(row_sum)) | (row_sum <= 0)
+    if bool(bad.any().item()):
+        fallback = torch.zeros_like(probs)
+        fallback.scatter_(-1, filtered_logits.argmax(dim=-1, keepdim=True), 1.0)
+        probs = torch.where(bad, fallback, probs)
+        row_sum = probs.sum(dim=-1, keepdim=True)
+    return probs / row_sum.clamp_min(1.0e-20)
+
+
+def _sample_top_p_sorted_tokens(logits, top_p):
+    """Sample from top-p filtered logits without scattering back to vocab order."""
+    sorted_logits, sorted_indices = torch.sort(logits, descending=True, dim=-1)
+    cumulative_probs = torch.cumsum(torch.softmax(sorted_logits, dim=-1), dim=-1)
+    remove = cumulative_probs > top_p
+    remove[..., 1:] = remove[..., :-1].clone()
+    remove[..., 0] = False
+    sorted_logits.masked_fill_(remove, torch.finfo(sorted_logits.dtype).min)
+    sorted_probs = _safe_probs(sorted_logits)
+    sample_idx = sorted_probs.argmax(dim=-1)
+    try:
+        sample_idx = torch.distributions.Categorical(probs=sorted_probs).sample()
+    except Exception:
+        pass
+    return sorted_indices.gather(-1, sample_idx.unsqueeze(-1)).squeeze(-1)
+
+
+@torch.no_grad()
+def sample_tokens_batched(logits, generated, token_ids, per_row_temp,
+                          repetition_penalty=1.0, top_p=None, top_k=None,
+                          keep_k_avg=4, generation_mode='fast'):
+    """Batched fork of generate_utils.sample_tokens for the MTP window [B,6,V]. The logits pipeline
+    (rep-penalty / per-row temperature / top_p / top_k / softmax / sample) is ROW-INDEPENDENT, so run
+    it ONCE over the whole batch instead of B times on [1,6,V] (the per-row san defeats batching by
+    slicing wlogits[b:b+1]). Only the variable-length box ASSEMBLY (decode_bbox_avg -> ragged shapes,
+    where sample_tokens' final torch.stack throws) stays per-row, returned as a LIST.
+
+    Equivalence to per-row san: every pipeline op reduces on dim=-1 only (never crosses the row dim),
+    so row b's processed logits/probs are bit-identical to slicing first -> greedy (per_row_temp==0,
+    argmax branch, no RNG) is BIT-EXACT. Under sampling, one batched Categorical changes the global
+    RNG consumption order vs B per-row draws -> box-size jitter (blessed; greedy is the exact gate).
+    apply_repetition_penalty already loops per-row internally, so passing the full [B,M] `generated`
+    is row-correct. keep_k_avg/generation_mode mirror sample_tokens' decode_bbox_avg call EXACTLY
+    (note: the per-row san passes keep_k=5 but decode_bbox_avg reads keep_k_avg, default 4 -- so 5 is
+    a no-op there; we replicate keep_k_avg=4). Returns (x0[B,6], boxes: list of B 1-D LongTensors)."""
+    gu = _gen_utils()
+    B, S, V = logits.shape                                        # S = N_FUTURE = 6
+    if repetition_penalty != 1.0:
+        logits = _apply_repetition_penalty_lowmem(logits, generated, repetition_penalty)
+    t = per_row_temp.to(dtype=logits.dtype).view(B, 1, 1)
+    sample_rows = per_row_temp > 0
+    if bool(sample_rows.all().item()):
+        logits.div_(t.clamp(min=1e-8))
+    elif bool(sample_rows.any().item()):
+        idx = sample_rows.nonzero(as_tuple=True)[0]
+        logits[idx].div_(t[idx].clamp(min=1e-8))
+    logits = _finite_logits_(logits)
+    if top_p is not None and top_p < 1:
+        logits = _top_p_logits_(logits, top_p)
+    if top_k is not None and top_k > 0:
+        logits = _top_k_logits_(logits, top_k)
+    probs = _safe_probs(logits)
+    x0 = probs.argmax(dim=-1)                                     # [B,6]; greedy rows are final here
+    samp = per_row_temp > 0
+    if bool(samp.any()):                                         # sampling rows: ONE batched Categorical draw
+        idx = samp.nonzero(as_tuple=True)[0]
+        try:
+            x0[idx] = gu.dists.Categorical(probs=probs[idx]).sample()
+        except Exception:
+            pass                                                # keep argmax (matches san's except: probs.max)
+    boxes = []
+    fallback = torch.zeros(1, dtype=x0.dtype, device=x0.device)
+    for b in range(B):                                          # variable-length box assembly (per-row, exact)
+        db = _decode_bbox_with_uncertainty(
+            logits[b], probs[b], token_ids,
+            keep_k=keep_k_avg, generation_mode=generation_mode)
+        if db is not None:
+            boxes.append(db)
+        else:
+            ref = gu.decode_ref(logits[b], probs[b], token_ids)
+            if ref is None:
+                boxes.append(fallback)
+            elif torch.is_tensor(ref):
+                boxes.append(ref.to(dtype=x0.dtype, device=x0.device))
+            else:
+                boxes.append(torch.tensor(ref, dtype=x0.dtype, device=x0.device))
+    return x0, boxes
+
+
+@torch.no_grad()
+def sample_next_tokens_batched(logits, generated, per_row_temp,
+                               repetition_penalty=1.0, top_p=None, top_k=None):
+    """Batched one-token sampler for AR repair rows.
+
+    This mirrors the row-independent part of ``sample_tokens`` for logits shaped
+    ``[B,1,V]``.  It intentionally does not run bbox/ref assembly because AR mode
+    only needs the next token before the state machine classifies it.
+    """
+    gu = _gen_utils()
+    if logits.dim() != 3 or logits.shape[1] != 1:
+        raise ValueError(f"AR batched sampler expects logits [B,1,V], got {tuple(logits.shape)}")
+    B = int(logits.shape[0])
+    if repetition_penalty != 1.0:
+        logits = _apply_repetition_penalty_lowmem(logits, generated, repetition_penalty)
+    t = per_row_temp.to(dtype=logits.dtype).view(B, 1, 1)
+    sample_rows = per_row_temp > 0
+    if bool(sample_rows.all().item()):
+        logits.div_(t.clamp(min=1e-8))
+    elif bool(sample_rows.any().item()):
+        idx = sample_rows.nonzero(as_tuple=True)[0]
+        logits[idx].div_(t[idx].clamp(min=1e-8))
+    logits = _finite_logits_(logits)
+    sorted_top_p = os.environ.get("AR_SORTED_TOPP", "0") == "1"
+    default_top_p = sorted_top_p and top_p is not None and top_p < 1 and (top_k is None or top_k <= 0)
+    if default_top_p and bool(sample_rows.all().item()):
+        return _sample_top_p_sorted_tokens(logits, top_p)
+    if top_p is not None and top_p < 1:
+        logits = _top_p_logits_(logits, top_p)
+    if top_k is not None and top_k > 0:
+        logits = _top_k_logits_(logits, top_k)
+    probs = _safe_probs(logits)
+    x0 = probs.argmax(dim=-1)
+    if bool(sample_rows.any().item()):
+        # Keep row-ordered sampling as the release default.  A single batched
+        # Categorical is faster, but it consumes RNG differently from stock AR
+        # repair and can alter default-temperature termination behavior.
+        for row in sample_rows.nonzero(as_tuple=True)[0].tolist():
+            try:
+                x0[row : row + 1] = gu.dists.Categorical(probs=probs[row : row + 1]).sample()
+            except Exception:
+                pass
+    return x0
+
+
+def load_pil(p):
+    from PIL import Image
+    im = Image.open(p).convert("RGB"); w, h = im.size
+    if max(w, h) > MAX_DIM:
+        s = MAX_DIM / max(w, h); im = im.resize((max(1, round(w*s)), max(1, round(h*s))), Image.LANCZOS)
+    return im
+
+def _preproc_one(im):
+    """CPU-side processor for one image -> (pixel_values[bf16], grid[int32]). Split out of
+    _encode_image so _encode_images can batch the GPU encode while preprocessing stays per-image."""
+    tok, proc, model = load()
+    msg = [{"role": "user", "content": [{"type": "image", "image": im}, {"type": "text", "text": "x"}]}]
+    text = proc.py_apply_chat_template(msg, tokenize=False, add_generation_prompt=True)
+    imgs, vids = proc.process_vision_info(msg)
+    inp = proc(text=[text], images=imgs, videos=vids, return_tensors="pt").to(DEV)
+    grid = inp.get("image_grid_hws")
+    if isinstance(grid, np.ndarray): grid = torch.from_numpy(grid).to(DEV, dtype=torch.int32)
+    return inp["pixel_values"].to(DT), grid
+
+
+def _vision_is_flash():
+    """True iff MoonViT will actually run flash_attn_varlen (so cross-image packing is
+    block-diagonal = exact AND a win). If the vision blocks are on sdpa/eager, OR the flash
+    wheel is absent (multihead_attention falls back to the dense-mask sdpa path), packing is
+    O(S^2) N^2 -> caller must stay per-image."""
+    vm = load()[2].vision_model
+    mod = importlib.import_module(type(vm).__module__)
+    if getattr(mod, "flash_attn_varlen_func", None) is None:
+        return False
+    try:
+        return vm.encoder.blocks[0].attn_implementation == "flash_attention_2"
+    except Exception:
+        return False
+
+
+@torch.no_grad()
+def _encode_images(ims):
+    """N images -> list of [n_img_tokens, C] mlp1-projected visual_features, one per image
+    (row-order). Drop-in for [_encode_image(im) for im in ims].
+
+    With flash present (_vision_is_flash) and N>1, packs images into
+    extract_feature micro-batches: MoonViT's varlen cu_seqlens path is
+    block-diagonal by image. Without flash, the dense SDPA fallback would scale
+    with the packed total sequence length, so this function falls back to
+    per-image encode. MTP_BATCH_VISION=0 also forces per-image encode."""
+    tok, proc, model = load()
+    pvs, grids = [], []
+    for im in ims:
+        pv, g = _preproc_one(im)
+        pvs.append(pv); grids.append(g)
+    if BATCH_VISION and len(ims) > 1 and _vision_is_flash():
+        if VISION_ENCODE_BATCH_SIZE <= 0 or VISION_ENCODE_BATCH_SIZE >= len(ims):
+            vit_list = model.extract_feature(torch.cat(pvs, dim=0), torch.cat(grids, dim=0))
+        else:
+            vit_list = []
+            for start in range(0, len(ims), VISION_ENCODE_BATCH_SIZE):
+                end = min(start + VISION_ENCODE_BATCH_SIZE, len(ims))
+                vit_list.extend(
+                    model.extract_feature(
+                        torch.cat(pvs[start:end], dim=0),
+                        torch.cat(grids[start:end], dim=0),
+                    )
+                )
+        return [model.mlp1(v) for v in vit_list]        # one [P_i, C] per image (patch_merger split)
+    return [model.mlp1(torch.cat(model.extract_feature(pv, g), dim=0))
+            for pv, g in zip(pvs, grids)]                # per-image (flash absent / N==1 / forced off)
+
+
+@torch.no_grad()
+def _encode_image(im):
+    """Single-image convenience wrapper (single-image callers); = _encode_images([im])[0]
+    (takes the per-image path inside _encode_images, so bit-identical to the original)."""
+    return _encode_images([im])[0]
+
+@torch.no_grad()
+def _tokenize(im, query):
+    """1-D prompt token ids for (image, query). Uses the model's own chat template."""
+    tok, proc, model = load()
+    msg = [{"role": "user", "content": [{"type": "image", "image": im},
+                                        {"type": "text", "text": _PROMPT + query + "."}]}]
+    text = proc.py_apply_chat_template(msg, tokenize=False, add_generation_prompt=True)
+    imgs, vids = proc.process_vision_info(msg)
+    return proc(text=[text], images=imgs, videos=vids, return_tensors="pt").to(DEV)["input_ids"][0]
+
+
+@torch.no_grad()
+def _tokenize_cached_image(query, image_token_count, im=None):
+    """Tokenize a prompt when the image token count is already known.
+
+    This keeps the processor's chat template, but directly expands ``<image-1>``
+    from the cached visual feature length.  It avoids re-running the CPU image
+    processor for every category prompt that shares the same image.
+    """
+    tok, proc, model = load()
+    msg = [{"role": "user", "content": [{"type": "image", "image": im},
+                                        {"type": "text", "text": _PROMPT + query + "."}]}]
+    text = proc.py_apply_chat_template(msg, tokenize=False, add_generation_prompt=True)
+    placeholder = f"<{getattr(proc, 'image_placeholder', 'image')}-1>"
+    image_token = getattr(proc, "image_token", "<IMG_CONTEXT>")
+    image_start = getattr(proc, "image_start_token", "<img>")
+    image_end = getattr(proc, "image_end_token", "</img>")
+    replacement = f"<image 1>{image_start}{image_token * int(image_token_count)}{image_end}"
+    if placeholder not in text:
+        raise ValueError(f"cached image placeholder {placeholder!r} was not found in chat template")
+    text = text.replace(placeholder, replacement, 1)
+    return tok([text], return_tensors="pt").to(DEV)["input_ids"][0]
+
+
+def _proc_full(im, query):
+    """Full processor dict (input_ids, attention_mask, pixel_values, image_grid_hws) —
+    used by the bench to drive the STOCK generate for the equivalence check."""
+    tok, proc, model = load()
+    msg = [{"role": "user", "content": [{"type": "image", "image": im},
+                                        {"type": "text", "text": _PROMPT + query + "."}]}]
+    text = proc.py_apply_chat_template(msg, tokenize=False, add_generation_prompt=True)
+    imgs, vids = proc.process_vision_info(msg)
+    inp = proc(text=[text], images=imgs, videos=vids, return_tensors="pt").to(DEV)
+    grid = inp.get("image_grid_hws")
+    if isinstance(grid, np.ndarray): grid = torch.from_numpy(grid).to(DEV, dtype=torch.int32)
+    inp["image_grid_hws"] = grid
+    return inp
+
+def _pad_generated(prompt_ids, gen_ids, img_tok, dev):
+    """Per-row [prompt + accepted] left-padded with the image token (already in every
+    prompt -> .unique() unchanged -> repetition penalty identical to single-run)."""
+    rows = [list(prompt_ids[b].tolist()) + gen_ids[b] for b in range(len(prompt_ids))]
+    M = max(len(r) for r in rows)
+    out = torch.full((len(rows), M), img_tok, dtype=torch.long, device=dev)
+    for b, r in enumerate(rows):
+        out[b, M - len(r):] = torch.tensor(r, dtype=torch.long, device=dev)
+    return out

chat_template.json

@@ -0,0 +1,4 @@
+{
+    "chat_template": "{% set image_count = namespace(value=0) %}{% set video_count = namespace(value=0) %}{% for message in messages %}{% if loop.first and message['role'] != 'system' %}<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n{% endif %}<|im_start|>{{ message['role'] }}\n{% if message['content'] is string %}{{ message['content'] }}<|im_end|>\n{% else %}{% for content in message['content'] %}{% if content['type'] == 'image' or 'image' in content or 'image_url' in content %}{% set image_count.value = image_count.value + 1 %}{% if add_vision_id %}<image {{ image_count.value }}>{% endif %}<image-{{ image_count.value }}>{% elif content['type'] == 'video' or 'video' in content %}{% set video_count.value = video_count.value + 1 %}{% if add_vision_id %}<video {{ video_count.value }}>{% endif %}<video-{{ video_count.value }}>{% elif 'text' in content %}{{ content['text'] }}{% endif %}{% endfor %}<|im_end|>\n{% endif %}{% endfor %}{% if add_generation_prompt %}<|im_start|>assistant\n{% endif %}"
+}
+  

config.json

@@ -0,0 +1,82 @@
+{
+  "_attn_implementation": "magi",
+  "_commit_hash": null,
+  "architectures": [
+    "LocateAnythingForConditionalGeneration"
+  ],
+  "auto_map": {
+    "AutoConfig": "configuration_locateanything.LocateAnythingConfig",
+    "AutoModel": "modeling_locateanything.LocateAnythingForConditionalGeneration"
+  },
+  "box_end_token_id": 151669,
+  "box_start_token_id": 151668,
+  "coord_end_token_id": 152677,
+  "coord_start_token_id": 151677,
+  "image_token_index": 151665,
+  "mlp_checkpoint": false,
+  "mlp_connector_layers": 2,
+  "model_type": "locateanything",
+  "none_token_id": 4064,
+  "ref_end_token_id": 151673,
+  "ref_start_token_id": 151672,
+  "template": null,
+  "text_config": {
+    "_attn_implementation_autoset": true,
+    "_name_or_path": "Qwen/Qwen2.5-3B-Instruct",
+    "architectures": [
+      "Qwen2ForCausalLM"
+    ],
+    "attention_dropout": 0.0,
+    "block_size": 6,
+    "bos_token_id": 151643,
+    "causal_attn": false,
+    "eos_token_id": 151645,
+    "hidden_act": "silu",
+    "hidden_size": 2048,
+    "initializer_range": 0.02,
+    "intermediate_size": 11008,
+    "max_position_embeddings": 32768,
+    "max_window_layers": 70,
+    "model_type": "qwen2",
+    "null_token_id": 152678,
+    "num_attention_heads": 16,
+    "num_hidden_layers": 36,
+    "num_key_value_heads": 2,
+    "rms_norm_eps": 1e-06,
+    "rope_scaling": null,
+    "rope_theta": 1000000.0,
+    "sliding_window": 32768,
+    "switch_token_id": 152679,
+    "text_mask_token_id": 151676,
+    "tie_word_embeddings": true,
+    "torch_dtype": "bfloat16",
+    "use_cache": false,
+    "use_sliding_window": false,
+    "vocab_size": 152681
+  },
+  "torch_dtype": "bfloat16",
+  "transformers_version": null,
+  "use_backbone_lora": 0,
+  "use_llm_lora": 0,
+  "vision_config": {
+    "_attn_implementation_autoset": true,
+    "_name_or_path": "moonshotai/MoonViT-SO-400M",
+    "auto_map": {
+      "AutoConfig": "moonshotai/MoonViT-SO-400M--configuration_moonvit.MoonViTConfig",
+      "AutoModel": "moonshotai/MoonViT-SO-400M--modeling_moonvit.MoonVitPretrainedModel"
+    },
+    "hidden_size": 1152,
+    "init_pos_emb_height": 64,
+    "init_pos_emb_width": 64,
+    "intermediate_size": 4304,
+    "merge_kernel_size": [
+      2,
+      2
+    ],
+    "model_type": "moonvit",
+    "num_attention_heads": 16,
+    "num_hidden_layers": 27,
+    "patch_size": 14,
+    "torch_dtype": "bfloat16"
+  }
+}

configuration_locateanything.py

@@ -0,0 +1,131 @@
+# --------------------------------------------------------
+# InternVL
+# Copyright (c) 2023 OpenGVLab
+# Licensed under The MIT License [see LICENSE for details]
+# --------------------------------------------------------
+
+import copy
+
+from transformers.models.qwen2.configuration_qwen2 import Qwen2Config
+from transformers.models.qwen3.configuration_qwen3 import Qwen3Config
+from transformers.configuration_utils import PretrainedConfig
+from transformers.utils import logging
+logger = logging.get_logger(__name__)
+
+class MoonViTConfig(PretrainedConfig):
+    model_type = "moonvit"
+
+    def __init__(
+        self,
+        patch_size: int = 14,
+        init_pos_emb_height: int = 64,
+        init_pos_emb_width: int = 64,
+        num_attention_heads: int = 16,
+        num_hidden_layers: int = 27,
+        hidden_size: int = 1152,
+        intermediate_size: int = 4304,
+        merge_kernel_size: tuple[int, int] = (2, 2),
+        **kwargs,
+    ):
+        super().__init__(**kwargs)
+        self.patch_size = patch_size
+        # Positional embedding config
+        self.init_pos_emb_height = init_pos_emb_height
+        self.init_pos_emb_width = init_pos_emb_width
+        # Transformer config
+        self.num_hidden_layers = num_hidden_layers
+        self.num_attention_heads = num_attention_heads
+        self.hidden_size = hidden_size
+        self.intermediate_size = intermediate_size
+        # Patch merger config
+        self.merge_kernel_size = merge_kernel_size
+
+
+class LocateAnythingConfig(PretrainedConfig):
+    model_type = 'locateanything'
+    is_composition = True
+    sub_configs = {"vision_config": MoonViTConfig, "text_config": Qwen2Config}
+    def __init__(
+            self,
+            vision_config=None,
+            text_config=None,
+            use_backbone_lora=0,
+            use_llm_lora=0,
+            downsample_ratio=0.5,
+            template=None,
+            loss_version='v1',
+            mlp_checkpoint=False,
+            image_token_index=151667,
+            box_start_token_id=151668,
+            box_end_token_id=151669,
+            coord_start_token_id=151677,
+            coord_end_token_id=152677,
+            ref_start_token_id=151672,
+            ref_end_token_id=151673,
+            none_token_id=4064,
+            **kwargs):
+        super().__init__(**kwargs)
+
+        if vision_config is None:
+            vision_config = {'model_type': 'moonvit'}
+            logger.info('vision_config is None. Initializing the MoonViTConfig with default values.')
+
+        if text_config is None:
+            text_config = {'architectures': ['Qwen2ForCausalLM']}
+            logger.info('text_config is None. Initializing the Qwen2Config config with default values.')
+
+        if vision_config['model_type'] == 'moonvit':
+            self.vision_config = MoonViTConfig(**vision_config)
+        else:
+            raise ValueError('Unsupported model_type: {}. Only moonvit is supported.'.format(vision_config['model_type']))
+
+
+        if text_config['architectures'][0] == 'Qwen2ForCausalLM':
+            self.text_config = Qwen2Config(**text_config)
+        elif text_config['architectures'][0] == 'Qwen3ForCausalLM':
+            self.text_config = Qwen3Config(**text_config)
+        else:
+            raise ValueError('Unsupported architecture: {}. Only Qwen2ForCausalLM and Qwen3ForCausalLM are supported.'.format(text_config['architectures'][0]))
+        self.use_backbone_lora = use_backbone_lora
+        self.use_llm_lora = use_llm_lora
+        self.mlp_checkpoint = mlp_checkpoint
+        self.downsample_ratio = downsample_ratio
+        self.template = template
+        self.loss_version = loss_version
+        self.tie_word_embeddings = self.text_config.tie_word_embeddings
+        self.image_token_index = image_token_index
+        self.box_start_token_id = box_start_token_id
+        self.box_end_token_id = box_end_token_id
+        self.coord_start_token_id = coord_start_token_id
+        self.coord_end_token_id = coord_end_token_id
+        self.ref_start_token_id = ref_start_token_id
+        self.ref_end_token_id = ref_end_token_id
+        self.none_token_id = none_token_id
+
+    def to_dict(self):
+        """
+        Serializes this instance to a Python dictionary. Override the default [`~PretrainedConfig.to_dict`].
+
+        Returns:
+            `Dict[str, any]`: Dictionary of all the attributes that make up this configuration instance,
+        """
+        output = copy.deepcopy(self.__dict__)
+        output['vision_config'] = self.vision_config.to_dict()
+        output['text_config'] = self.text_config.to_dict()
+        output['model_type'] = self.__class__.model_type
+        output['use_backbone_lora'] = self.use_backbone_lora
+        output['use_llm_lora'] = self.use_llm_lora
+        output['downsample_ratio'] = self.downsample_ratio
+        output['template'] = self.template
+        output['image_token_index'] = self.image_token_index
+        output['box_start_token_id'] = self.box_start_token_id
+        output['box_end_token_id'] = self.box_end_token_id
+        output['coord_start_token_id'] = self.coord_start_token_id
+        output['coord_end_token_id'] = self.coord_end_token_id
+        output['ref_start_token_id'] = self.ref_start_token_id
+        output['ref_end_token_id'] = self.ref_end_token_id
+        output['none_token_id'] = self.none_token_id
+        output['_attn_implementation'] = self._attn_implementation
+        if hasattr(self, '_attn_implementation_autoset'):
+            output['_attn_implementation_autoset'] = self._attn_implementation_autoset
+        return output

configuration_qwen2.py

@@ -0,0 +1,148 @@
+# coding=utf-8
+# Copyright 2024 The Qwen team, Alibaba Group 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.
+""" Qwen2 model configuration"""
+
+from transformers.configuration_utils import PretrainedConfig
+from transformers.utils import logging
+
+
+logger = logging.get_logger(__name__)
+
+QWEN2_PRETRAINED_CONFIG_ARCHIVE_MAP = {
+    "Qwen/Qwen2-7B-beta": "https://huggingface.co/Qwen/Qwen2-7B-beta/resolve/main/config.json",
+}
+
+
+class Qwen2Config(PretrainedConfig):
+    r"""
+    This is the configuration class to store the configuration of a [`Qwen2Model`]. It is used to instantiate a
+    Qwen2 model according to the specified arguments, defining the model architecture. Instantiating a configuration
+    with the defaults will yield a similar configuration to that of
+    Qwen2-7B-beta [Qwen/Qwen2-7B-beta](https://huggingface.co/Qwen/Qwen2-7B-beta).
+
+    Configuration objects inherit from [`PretrainedConfig`] and can be used to control the model outputs. Read the
+    documentation from [`PretrainedConfig`] for more information.
+
+
+    Args:
+        vocab_size (`int`, *optional*, defaults to 151936):
+            Vocabulary size of the Qwen2 model. Defines the number of different tokens that can be represented by the
+            `inputs_ids` passed when calling [`Qwen2Model`]
+        hidden_size (`int`, *optional*, defaults to 4096):
+            Dimension of the hidden representations.
+        intermediate_size (`int`, *optional*, defaults to 22016):
+            Dimension of the MLP representations.
+        num_hidden_layers (`int`, *optional*, defaults to 32):
+            Number of hidden layers in the Transformer encoder.
+        num_attention_heads (`int`, *optional*, defaults to 32):
+            Number of attention heads for each attention layer in the Transformer encoder.
+        num_key_value_heads (`int`, *optional*, defaults to 32):
+            This is the number of key_value heads that should be used to implement Grouped Query Attention. If
+            `num_key_value_heads=num_attention_heads`, the model will use Multi Head Attention (MHA), if
+            `num_key_value_heads=1 the model will use Multi Query Attention (MQA) otherwise GQA is used. When
+            converting a multi-head checkpoint to a GQA checkpoint, each group key and value head should be constructed
+            by meanpooling all the original heads within that group. For more details checkout [this
+            paper](https://arxiv.org/pdf/2305.13245.pdf). If it is not specified, will default to `32`.
+        hidden_act (`str` or `function`, *optional*, defaults to `"silu"`):
+            The non-linear activation function (function or string) in the decoder.
+        max_position_embeddings (`int`, *optional*, defaults to 32768):
+            The maximum sequence length that this model might ever be used with.
+        initializer_range (`float`, *optional*, defaults to 0.02):
+            The standard deviation of the truncated_normal_initializer for initializing all weight matrices.
+        rms_norm_eps (`float`, *optional*, defaults to 1e-06):
+            The epsilon used by the rms normalization layers.
+        use_cache (`bool`, *optional*, defaults to `True`):
+            Whether or not the model should return the last key/values attentions (not used by all models). Only
+            relevant if `config.is_decoder=True`.
+        tie_word_embeddings (`bool`, *optional*, defaults to `False`):
+            Whether the model's input and output word embeddings should be tied.
+        rope_theta (`float`, *optional*, defaults to 10000.0):
+            The base period of the RoPE embeddings.
+        use_sliding_window (`bool`, *optional*, defaults to `False`):
+            Whether to use sliding window attention.
+        sliding_window (`int`, *optional*, defaults to 4096):
+            Sliding window attention (SWA) window size. If not specified, will default to `4096`.
+        max_window_layers (`int`, *optional*, defaults to 28):
+            The number of layers that use SWA (Sliding Window Attention). The bottom layers use SWA while the top use full attention.
+        attention_dropout (`float`, *optional*, defaults to 0.0):
+            The dropout ratio for the attention probabilities.
+
+    ```python
+    >>> from transformers import Qwen2Model, Qwen2Config
+
+    >>> # Initializing a Qwen2 style configuration
+    >>> configuration = Qwen2Config()
+
+    >>> # Initializing a model from the Qwen2-7B style configuration
+    >>> model = Qwen2Model(configuration)
+
+    >>> # Accessing the model configuration
+    >>> configuration = model.config
+    ```"""
+
+    model_type = "qwen2"
+    keys_to_ignore_at_inference = ["past_key_values"]
+
+    def __init__(
+        self,
+        vocab_size=151936,
+        hidden_size=4096,
+        intermediate_size=22016,
+        num_hidden_layers=32,
+        num_attention_heads=32,
+        num_key_value_heads=32,
+        hidden_act="silu",
+        max_position_embeddings=32768,
+        initializer_range=0.02,
+        rms_norm_eps=1e-6,
+        use_cache=True,
+        tie_word_embeddings=False,
+        rope_theta=10000.0,
+        use_sliding_window=False,
+        sliding_window=4096,
+        max_window_layers=28,
+        attention_dropout=0.0,
+        **kwargs,
+    ):
+        self.vocab_size = vocab_size
+        self.max_position_embeddings = max_position_embeddings
+        self.hidden_size = hidden_size
+        self.intermediate_size = intermediate_size
+        self.num_hidden_layers = num_hidden_layers
+        self.num_attention_heads = num_attention_heads
+        self.use_sliding_window = use_sliding_window
+        self.sliding_window = sliding_window
+        self.max_window_layers = max_window_layers
+
+        # for backward compatibility
+        if num_key_value_heads is None:
+            num_key_value_heads = num_attention_heads
+
+        self.num_key_value_heads = num_key_value_heads
+        self.hidden_act = hidden_act
+        self.initializer_range = initializer_range
+        self.rms_norm_eps = rms_norm_eps
+        self.use_cache = use_cache
+        self.rope_theta = rope_theta
+        self.attention_dropout = attention_dropout
+        if kwargs.get('attn_implementation', None) is None:
+            self.attn_implementation = kwargs['attn_implementation'] = 'flash_attention_2'
+        else:
+            self.attn_implementation = kwargs['attn_implementation']
+
+        super().__init__(
+            tie_word_embeddings=tie_word_embeddings,
+            **kwargs,
+        )

generate_utils.py

@@ -0,0 +1,504 @@
+# Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
+#
+# NVIDIA CORPORATION and its licensors retain all intellectual property
+# and proprietary rights in and to this software, related documentation
+# and any modifications thereto.  Any use, reproduction, disclosure or
+# distribution of this software and related documentation without an express
+# license agreement from NVIDIA CORPORATION is strictly prohibited.
+
+import torch
+import torch.nn.functional as F
+import torch.distributions as dists
+from typing import Dict, Optional
+
+
+def get_token_ids_from_config(config) -> Dict[str, int]:
+    """Extract all token IDs from the configuration object.
+    
+    Args:
+        config: Configuration object (LocateAnythingConfig or similar)
+        
+    Returns:
+        Dictionary containing all token IDs
+    """
+    token_ids = {}
+    
+    # Get from main config
+    token_ids['box_start_token_id'] = getattr(config, 'box_start_token_id', 151668)
+    token_ids['box_end_token_id'] = getattr(config, 'box_end_token_id', 151669)
+    token_ids['coord_start_token_id'] = getattr(config, 'coord_start_token_id', 151677)
+    token_ids['coord_end_token_id'] = getattr(config, 'coord_end_token_id', 152677)
+    token_ids['ref_start_token_id'] = getattr(config, 'ref_start_token_id', 151672)
+    token_ids['ref_end_token_id'] = getattr(config, 'ref_end_token_id', 151673)
+    token_ids['none_token_id'] = getattr(config, 'none_token_id', 4064)
+    
+    # Get from text_config
+    text_config = getattr(config, 'text_config', None)
+    if text_config is not None:
+        token_ids['null_token_id'] = getattr(text_config, 'null_token_id', 152678)
+        token_ids['im_end_token_id'] = getattr(text_config, 'eos_token_id', 151645)
+        token_ids['switch_token_id'] = getattr(text_config, 'switch_token_id', 152679)
+        token_ids['default_mask_token_id'] = getattr(text_config, 'text_mask_token_id', 151676)
+    else:
+        token_ids['null_token_id'] = 152678
+        token_ids['im_end_token_id'] = 151645
+        token_ids['switch_token_id'] = 152679
+        token_ids['default_mask_token_id'] = 151676
+    
+    return token_ids
+
+
+def top_p_logits(
+    logits: torch.Tensor,
+    top_p: float = None
+) -> torch.Tensor:
+    sorted_logits, sorted_indices = torch.sort(logits, descending=True)
+    cumulative_probs = torch.cumsum(F.softmax(sorted_logits, dim=-1), dim=-1)
+    sorted_indices_to_remove = cumulative_probs > top_p
+    # Shift the indices to the right to keep the first token above the threshold
+    sorted_indices_to_remove[..., 1:] = sorted_indices_to_remove[..., :-1].clone()
+    sorted_indices_to_remove[..., 0] = 0
+
+    mask = torch.zeros_like(logits, dtype=torch.bool, device=logits.device)
+    mask = mask.scatter_(-1, sorted_indices, sorted_indices_to_remove)
+    logits = logits.masked_fill(mask, torch.finfo(logits.dtype).min)
+    return logits
+
+
+def top_k_logits(
+    logits: torch.Tensor,
+    top_k: int = None
+) -> torch.Tensor:
+    top_k = min(top_k, logits.size(-1))  # Safety check
+    # Remove all tokens with a probability less than the last token of the top-k
+    indices_to_remove = logits < torch.topk(logits, top_k)[0][..., -1, None]
+    logits = logits.masked_fill(indices_to_remove, torch.finfo(logits.dtype).min)
+    return logits
+
+
+def apply_repetition_penalty(
+    logits: torch.Tensor,
+    input_ids: torch.Tensor,
+    repetition_penalty: float = 1.0
+) -> torch.Tensor:
+    """
+    Apply repetition penalty to logits.
+    
+    Args:
+        logits: Shape [batch_size, seq_len, vocab_size] or [batch_size, vocab_size]
+        input_ids: Previously generated token ids, shape [batch_size, seq_len]
+        repetition_penalty: Penalty factor. > 1.0 penalizes repetition, < 1.0 encourages it.
+    
+    Returns:
+        Modified logits with repetition penalty applied.
+    """
+    if repetition_penalty == 1.0:
+        return logits
+    
+    # Convert to 3D for vectorized computation
+    if logits.dim() == 2:
+        logits = logits.unsqueeze(1)  # [B, 1, V]
+        squeeze_back = True
+    else:
+        squeeze_back = False
+    
+    batch_size, seq_len, vocab_size = logits.shape
+
+    # Construct [B, V] bool mask marking tokens that have appeared in each batch
+    device = logits.device
+    token_mask = torch.zeros(batch_size, vocab_size, dtype=torch.bool, device=device)
+    for b in range(batch_size):
+        # Apply penalty only based on tokens already generated in this batch
+        unique_tokens = input_ids[b].unique()
+        # Prevent out-of-bounds: only keep IDs within vocab range
+        valid_tokens = unique_tokens[(unique_tokens >= 0) & (unique_tokens < vocab_size)]
+        if valid_tokens.numel() > 0:
+            token_mask[b, valid_tokens] = True
+
+    # Expand to [B, L, V] to align with logits
+    token_mask = token_mask.unsqueeze(1).expand(-1, seq_len, -1)
+
+    # Divide positive values by penalty, multiply negative values by penalty
+    positive = logits > 0
+    negative = ~positive
+
+    # Apply penalty only at mask positions
+    logits = torch.where(token_mask & positive, logits / repetition_penalty, logits)
+    logits = torch.where(token_mask & negative, logits * repetition_penalty, logits)
+    
+    if squeeze_back:
+        logits = logits.squeeze(1)
+    
+    return logits
+
+
+def sample_tokens(
+    logits: torch.Tensor,
+    generated: torch.Tensor,
+    token_ids: Dict[str, int],
+    **generate_kwargs,
+):  
+    batch_size, seq_len, vocab_size = logits.shape
+
+    repetition_penalty = generate_kwargs.get('repetition_penalty', 1.0)
+    temperature = generate_kwargs.get('temperature', 0)
+    top_p = generate_kwargs.get('top_p', None)
+    top_k = generate_kwargs.get('top_k', None)
+
+    # Apply repetition penalty based on all previously generated tokens
+    if repetition_penalty != 1.0:
+        logits = apply_repetition_penalty(logits, generated, repetition_penalty)
+
+    if temperature > 0:
+            logits = logits / temperature
+    if top_p is not None and top_p < 1:
+        logits = top_p_logits(logits, top_p)
+    if top_k is not None:
+        logits = top_k_logits(logits, top_k)
+    
+    probs = torch.softmax(logits, dim=-1)
+
+    if temperature > 0:
+        try:
+            x0 = dists.Categorical(probs=probs).sample()
+            confidence = torch.gather(probs, -1, x0.unsqueeze(-1)).squeeze(-1)
+        except Exception:
+            confidence, x0 = probs.max(dim=-1)
+    else:
+        confidence, x0 = probs.max(dim=-1)
+
+    if seq_len == 1:
+        return probs, confidence, x0, None
+
+    box_avg = []
+    fallback_box = torch.zeros(1, dtype=x0.dtype, device=x0.device)
+
+    for b in range(batch_size):
+        decoded_box = decode_bbox_avg(
+            logits[b], probs[b], token_ids, keep_k=generate_kwargs.get('keep_k_avg', 4),
+            generation_mode=generate_kwargs.get('generation_mode', 'hybrid'),
+        )
+        if decoded_box is not None:
+            box_avg.append(decoded_box)
+        else:
+            out_ref = decode_ref(logits[b], probs[b], token_ids)
+            if out_ref is not None:
+                box_avg.append(torch.tensor(out_ref, dtype=x0.dtype, device=x0.device))
+            else:
+                box_avg.append(fallback_box)
+
+    box_avg = torch.stack(box_avg)
+
+    return probs, confidence, x0, box_avg
+
+
+def sample_tokens_ar(
+    logits: torch.Tensor,
+    generated: torch.Tensor,
+    token_ids: Dict[str, int],
+    **generate_kwargs,
+):
+    """
+    Lightweight sampling function for AR single-step sampling only.
+    
+    Args:
+        logits: [batch_size, vocab_size] or [batch_size, 1, vocab_size]
+        generated: [batch_size, seq_len]
+    """
+    # Convert to 3D for reusing repetition penalty and clipping logic
+    if logits.dim() == 2:
+        logits = logits.unsqueeze(1)  # [B, 1, V]
+    batch_size, seq_len, vocab_size = logits.shape
+    assert seq_len == 1, "sample_tokens_ar only supports single-step AR sampling (seq_len == 1)"
+
+    repetition_penalty = generate_kwargs.get('repetition_penalty', 1.0)
+    temperature = generate_kwargs.get('temperature', 0)
+    top_p = generate_kwargs.get('top_p', None)
+    top_k = generate_kwargs.get('top_k', None)
+
+    # Apply repetition penalty only based on historically generated tokens
+    if repetition_penalty != 1.0:
+        logits = apply_repetition_penalty(logits, generated, repetition_penalty)
+
+    if temperature > 0:
+        logits = logits / temperature
+    if top_p is not None and top_p < 1:
+        logits = top_p_logits(logits, top_p)
+    if top_k is not None:
+        logits = top_k_logits(logits, top_k)
+
+    probs = torch.softmax(logits, dim=-1)
+
+    if temperature > 0:
+        try:
+            x0 = dists.Categorical(probs=probs).sample()
+            confidence = torch.gather(probs, -1, x0.unsqueeze(-1)).squeeze(-1)
+        except Exception:
+            confidence, x0 = probs.max(dim=-1)
+    else:
+        # For greedy: directly take the token with maximum probability
+        confidence, x0 = probs.max(dim=-1)
+
+    # Keep interface consistent with sample_tokens: return [B, 1, V] / [B, 1] shape
+    return probs, confidence, x0, None, None
+
+
+def is_valid_box_frame(
+    probs,
+    token_ids: Dict[str, int],
+    start_thresh=0.6,
+    end_thresh=0.2,
+    topk=5,
+):
+    box_start_token_id = token_ids['box_start_token_id']
+    box_end_token_id = token_ids['box_end_token_id']
+    null_token_id = token_ids['null_token_id']
+    im_end_token_id = token_ids['im_end_token_id']
+    none_token_id = token_ids['none_token_id'] # none
+
+    p_start = probs[0, box_start_token_id]
+    if p_start >= start_thresh:
+        if (probs[1, none_token_id] > 0.2 and 
+            probs[2, box_end_token_id] > 0.2 and 
+            probs[3, null_token_id] > 0.1 and 
+            probs[4, null_token_id] > 0.1):
+            return 'empty_box'
+
+    end_target_ids = torch.tensor([box_end_token_id, null_token_id, im_end_token_id], device=probs.device)
+    end_score = probs[5, end_target_ids].sum()
+
+    if end_score >= end_thresh:
+            return 'legal_box'
+
+    return 'illegal_box'
+
+
+def decode_bbox_avg(
+    logits,
+    probs,
+    token_ids: Dict[str, int],
+    keep_k=5,
+    start_thresh=0.7,
+    end_thresh=0.2,
+    generation_mode: str = 'hybrid',
+):
+    """
+    Decode bounding box coordinates using top-k weighted average.
+    
+    Args:
+        logits: Logits of shape (6, vocab_size)
+        probs: Probability distribution of shape (6, vocab_size)
+        token_ids: Dictionary containing all token IDs
+        keep_k: Number of top-k candidate tokens to keep at each position
+        start_thresh: Confidence threshold for box start token
+        end_thresh: Confidence threshold for box end token
+        
+    Returns:
+        Decoded bounding box coordinate list in format [box_start, x1, x2, y1, y2, box_end],
+        or None if decoding fails
+    """
+    coord_start_token_id = token_ids['coord_start_token_id']
+    coord_end_token_id = token_ids['coord_end_token_id']
+    box_start_token_id = token_ids['box_start_token_id']
+    box_end_token_id = token_ids['box_end_token_id']
+    none_token_id = token_ids['none_token_id']
+
+    device = logits.device
+
+    box_type = is_valid_box_frame(
+        probs,
+        token_ids,
+        start_thresh=start_thresh,
+        end_thresh=end_thresh,
+        topk=keep_k
+    )
+    if box_type == 'empty_box':
+        # Handle the <box>none</box> case first
+        return torch.tensor([
+            box_start_token_id, 
+            none_token_id, 
+            box_end_token_id, 
+            token_ids['null_token_id'], 
+            token_ids['null_token_id'], 
+            token_ids['null_token_id']
+        ], dtype=torch.long, device=probs.device)
+    elif box_type == 'illegal_box':
+        return None
+
+    # Extract probabilities at positions 1-4 and compute Top-K for all 4 positions at once
+    pos_probs, pos_ids = torch.topk(probs[1:5], k=keep_k, dim=-1)
+    mask = (pos_ids >= coord_start_token_id) & (pos_ids <= coord_end_token_id)
+    has_valid = mask.any(dim=-1) # shape: [4]
+    if not has_valid.all():
+        return None # not a box, exit...
+
+    first_valid_idx = mask.long().argmax(dim=-1, keepdim=True) # [4, 1]
+    # Extract highest-probability valid_probs[0] and corresponding valid_ids[0]
+    first_valid_probs = pos_probs.gather(-1, first_valid_idx).squeeze(-1) # [4]
+    first_valid_ids = pos_ids.gather(-1, first_valid_idx).squeeze(-1) # [4]
+    if generation_mode == 'hybrid':
+        valid_counts = mask.sum(dim=-1) # [4]
+        # Compute max/min of valid ids: fill invalid positions with extreme values to avoid interfering with max/min
+        LARGE_NUM, SMALL_NUM = 999999, -999999
+        valid_ids_for_max = torch.where(mask, pos_ids, torch.tensor(SMALL_NUM, device=device))
+        valid_ids_for_min = torch.where(mask, pos_ids, torch.tensor(LARGE_NUM, device=device))
+
+        valid_max = valid_ids_for_max.max(dim=-1)[0]
+        valid_min = valid_ids_for_min.min(dim=-1)[0]
+
+        is_abnormal = (first_valid_probs < 0.9) & (valid_counts > 1) & ((valid_max - valid_min) > 60)
+        # is_abnormal = (first_valid_probs < 0.7) & (valid_counts > 1) & ((valid_max - valid_min) > 80)
+
+        # Normal positions take top-1 (first_valid_ids); abnormal positions are replaced with 0
+        final_coords = torch.where(is_abnormal, torch.tensor(0, device=pos_ids.device), first_valid_ids)
+    elif generation_mode == 'fast':
+        final_coords = first_valid_ids
+
+
+    start_t = torch.tensor([box_start_token_id], dtype=final_coords.dtype, device=device)
+    end_t = torch.tensor([box_end_token_id], dtype=final_coords.dtype, device=device)
+
+    return torch.cat([start_t, final_coords, end_t])
+    
+
+def decode_ref(
+    logits,
+    probs,
+    token_ids: Dict[str, int],
+    keep_k=5,
+    start_thresh=0.6,
+):
+    ref_start_token_id = token_ids.get('ref_start_token_id')
+    coord_start_token_id = token_ids['coord_start_token_id']
+    coord_end_token_id = token_ids['coord_end_token_id']
+    device = probs.device
+    L = probs.size(0)
+
+    # 1. Check if the first position is <ref> and its probability meets start_thresh
+    # Note: we directly use the probability of the ref token at position 0 for the check
+    if probs[0, ref_start_token_id] < start_thresh:
+        return None
+    
+    # 2. Extract Top-K probabilities and token IDs for all subsequent positions
+    pos_probs, pos_ids = torch.topk(probs[1:], k=keep_k, dim=-1) # shape: [L-1, keep_k]
+    
+    # 3. Build mask: identify coordinate tokens (<0> ~ <1000>)
+    is_coord = (pos_ids >= coord_start_token_id) & (pos_ids <= coord_end_token_id)
+    # Invert: valid tokens are non-coordinate tokens
+    is_valid = ~is_coord # shape: [L-1, keep_k]
+
+    # Ensure each position has at least one non-coordinate valid token in its Top-K
+    has_valid = is_valid.any(dim=-1) # shape: [L-1]
+    if not has_valid.all():
+        return None
+
+    # 4. Get the highest-probability valid token
+    # Since topk results are sorted in descending order of probability,
+    # argmax returns the first index where is_valid is True, i.e., the index of the most probable valid token
+    first_valid_idx = is_valid.long().argmax(dim=-1, keepdim=True) # shape: [L-1, 1]
+    
+    # Extract the final token IDs
+    final_text_ids = pos_ids.gather(-1, first_valid_idx).squeeze(-1) # shape: [L-1]
+
+    start_t = torch.tensor([ref_start_token_id], dtype=final_text_ids.dtype, device=device)
+    
+    return torch.cat([start_t, final_text_ids])
+
+
+def handle_pattern(x0, token_ids: Dict[str, int], generation_mode: str = 'hybrid'):
+    """
+    Args:
+        x0: Token ID list of length 6
+        token_ids: Dictionary containing all token IDs
+    """
+    null_token_id = token_ids['null_token_id']
+    im_end_token_id = token_ids['im_end_token_id']
+    box_start_token_id = token_ids['box_start_token_id']
+    box_end_token_id = token_ids['box_end_token_id']
+    none_token_id = token_ids['none_token_id']
+    coord_start_token_id = token_ids['coord_start_token_id']
+    coord_end_token_id = token_ids['coord_end_token_id']
+    ref_end_token_id = token_ids['ref_end_token_id']
+    
+    x0 = x0.tolist()
+
+    if x0[0] == null_token_id:
+        return {
+            "type": "im_end",
+            "tokens": [im_end_token_id],
+            "need_switch_to_ar": False,
+            "is_terminal": True,
+        }
+    elif x0[0] == im_end_token_id:
+        return {
+            "type": "im_end",
+            "tokens": [im_end_token_id],
+            "need_switch_to_ar": False,
+            "is_terminal": True,
+        }
+    elif x0[:2] == [box_start_token_id, none_token_id]:
+        return {
+            "type": "empty_box",
+            "tokens": [box_start_token_id, none_token_id, box_end_token_id],
+            "need_switch_to_ar": False,
+            "is_terminal": False,
+        }
+    elif x0[0] == box_start_token_id:
+        coord_ix = 1
+        for coord in x0[1:5]:
+            if coord_start_token_id <= coord <= coord_end_token_id:
+                coord_ix += 1
+            else:
+                break
+
+        # Standard 4-coordinate bbox: <box><x1><x2><y1><y2></box>
+        if coord_ix == 5 and x0[5] == box_end_token_id:
+            return {
+                "type": "coord_box",
+                "tokens": x0,
+                "need_switch_to_ar": False,
+                "is_terminal": False,
+            }
+        # Two-coordinate pointing: <box><x><y></box>
+        # Convention: the first two coordinates are valid coord tokens, the third token is box_end.
+        # Remaining positions (if any) are not part of the pattern; truncate at box_end.
+        elif coord_ix == 3 and x0[3] == box_end_token_id:
+            return {
+                "type": "point_box",
+                "tokens": x0[:4],
+                "need_switch_to_ar": False,
+                "is_terminal": False,
+            }
+        else:
+            if generation_mode == 'fast':
+                # fast mode: treat as coord_box, stay in MTP
+                return {
+                    "type": "coord_box",
+                    "tokens": x0,
+                    "need_switch_to_ar": False,
+                    "is_terminal": False,
+                }
+            else:
+                # hybrid mode: error_box, switch to AR
+                return {
+                    "type": "error_box",
+                    "tokens": x0[:coord_ix],
+                    "need_switch_to_ar": True,
+                    "is_terminal": False,
+                }
+
+    else:
+        for i, token in enumerate(x0):
+            if token == null_token_id:
+                x0 = x0[:i]
+                break
+
+        if len(x0) >= 2 and x0[-1] == x0[-2] == ref_end_token_id:
+            x0 = x0[:-1]
+
+        return {
+            "type": "ref_object",
+            "tokens": x0,
+            "need_switch_to_ar": False,
+            "is_terminal": False,
+        }

generation_config.json

@@ -0,0 +1,6 @@
+{
+  "_from_model_config": true,
+  "bos_token_id": 151643,
+  "eos_token_id": 151645,
+  "transformers_version": "4.51.0"
+}

image_processing_locateanything.py

@@ -0,0 +1,128 @@
+# Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
+#
+# NVIDIA CORPORATION and its licensors retain all intellectual property
+# and proprietary rights in and to this software, related documentation
+# and any modifications thereto.  Any use, reproduction, disclosure or
+# distribution of this software and related documentation without an express
+# license agreement from NVIDIA CORPORATION is strictly prohibited.
+
+"""Image processor class for KimiVL."""
+
+import math
+import numpy as np
+from PIL import Image
+from typing import Optional, Union
+
+import torch
+from torchvision.transforms import functional as TF
+from transformers.image_utils import ImageInput, make_list_of_images, valid_images
+from transformers.image_processing_utils import BaseImageProcessor, BatchFeature
+from transformers.utils import TensorType
+from transformers import AutoImageProcessor
+
+MEAN = (0.5, 0.5, 0.5)
+STD = (0.5, 0.5, 0.5)
+
+
+class LocateAnythingImageProcessor(BaseImageProcessor):
+    model_type = "locateanything"
+
+    def __init__(
+        self,
+        patch_size: int = 14,
+        image_mean: tuple[float, float, float] = MEAN,
+        image_std: tuple[float, float, float] = STD,
+        in_token_limit: int = 4096,
+        merge_kernel_size: list[int, int] = [2, 2],
+        **kwargs,
+    ):
+        super().__init__(**kwargs)
+        self.in_token_limit = in_token_limit
+        self.patch_size = patch_size
+        self.image_mean = image_mean
+        self.image_std = image_std
+        self.merge_kernel_size = merge_kernel_size
+
+    def rescale(
+        self, image: Image.Image, merge_kernel_size: list[int, int] = [2, 2]
+    ) -> Image.Image:
+        w, h = image.size
+        patch_size = self.patch_size
+
+        if (w // patch_size) * (h // patch_size) > self.in_token_limit:
+            scale = math.sqrt(self.in_token_limit / ((w // patch_size) * (h // patch_size)))
+            new_w, new_h = int(w * scale), int(h * scale)
+            image = image.resize((new_w, new_h), Image.Resampling.BICUBIC)
+
+        new_w, new_h = image.size
+        pad_size_h = merge_kernel_size[0] * patch_size
+        pad_size_w = merge_kernel_size[1] * patch_size
+
+        target_w = math.ceil(new_w / pad_size_w) * pad_size_w
+        target_h = math.ceil(new_h / pad_size_h) * pad_size_h
+
+        if target_w != new_w or target_h != new_h:
+            image = image.resize((target_w, target_h), Image.Resampling.BICUBIC)
+
+        w, h = image.size
+        if w // patch_size >= 512 or h // patch_size >= 512:
+            raise ValueError("Exceed pos emb")
+
+        return image
+
+    def to_tensor(self, image: Image.Image) -> torch.Tensor:
+        return TF.to_tensor(image.convert("RGB"))
+
+    def normalize(self, image: torch.Tensor) -> torch.Tensor:
+        return TF.normalize(image, self.image_mean, self.image_std)
+
+    def patchify(self, image: torch.Tensor) -> tuple[torch.Tensor, list[int, int]]:
+        patch_size = self.patch_size
+        C, H, W = image.shape
+        patches = image.reshape(C, H // patch_size, patch_size, W // patch_size, patch_size)
+        patches = patches.permute(1, 3, 0, 2, 4)
+        patches = patches.contiguous().view(-1, C, patch_size, patch_size)
+        grid_hw = (H // patch_size, W // patch_size)
+        return patches, grid_hw
+
+    def _preprocess(self, image: ImageInput) -> tuple[torch.Tensor, list[int, int]]:
+        """
+        Preprocess image and patchify it.
+        Args:
+            image (`ImageInput`):
+                Image to preprocess. Expects pixel values ranging from 0 to 255. If pixel values range from 0 to 1, set `do_rescale=False`.
+        Returns:
+            patches: torch.Tensor
+            grid_hw: list[int, int]
+        """
+        image = self.rescale(image, self.merge_kernel_size)
+        image = self.to_tensor(image)
+        image = self.normalize(image)
+        patches, grid_hw = self.patchify(image)
+        return patches, grid_hw
+
+    def preprocess(
+        self,
+        images: ImageInput,
+        return_tensors: Optional[Union[str, TensorType]] = None,
+    ) -> BatchFeature:
+        images = make_list_of_images(images)
+
+        if not valid_images(images):
+            raise ValueError(
+                "Invalid image type. Must be of type PIL.Image.Image, numpy.ndarray, "
+                "torch.Tensor, tf.Tensor or jax.ndarray."
+            )
+
+        pixel_values, image_grid_hws = [], []
+        for image in images:
+            patches, image_grid_hw = self._preprocess(image)
+            pixel_values.append(patches)
+            image_grid_hws.append(image_grid_hw)
+        pixel_values = torch.concat(pixel_values, dim=0)
+        image_grid_hws = np.array(image_grid_hws)
+        data = {"pixel_values": pixel_values, "image_grid_hws": image_grid_hws}
+
+        return BatchFeature(data=data, tensor_type=return_tensors)
+
+AutoImageProcessor.register("LocateAnythingImageProcessor", LocateAnythingImageProcessor)

kernel_utils/README.md

@@ -0,0 +1,75 @@
+# LA Flash Utils
+
+This folder contains the sparse attention utilities used by
+`LA_FLASH_ATTN=la_flash`. The release path is implemented with
+FlashAttention varlen over LocateAnything range plans. It does not include or
+build a local C++/CUDA extension.
+
+## Features
+
+- Supports batched LocateAnything hybrid MTP inference on A100, RTX 4090, and H100.
+- Consumes Magi-style `q_ranges`, `k_ranges`, `segment_offsets`, and
+  `attn_type_map` plans generated by `batch_utils.hybrid_runtime`.
+- Uses FlashAttention varlen for packed causal/full plans.
+- Packs LocateAnything MTP full-window key segments before calling
+  FlashAttention, avoiding dense `[B,H,Q,K]` masks.
+- Supports log-sum-exp merging for compatible non-packed multi-segment plans.
+
+## Attention Types
+
+The release path intentionally supports only FlashAttention-compatible plan
+types:
+
+| Value | Meaning |
+| --- | --- |
+| `0` | Full attention over the listed key segment or packed key segments. |
+| `1` | Bottom-right causal attention. |
+
+## How It Works
+
+`batch_utils.hybrid_runtime` builds sparse range plans for the text decoder.
+Each plan describes which query token intervals attend to which key/value token
+intervals. `kernel_utils.range_attention` executes those plans with
+FlashAttention instead of materializing dense SDPA masks.
+
+The runtime follows three paths:
+
+- **Packed simple plans:** when each query range maps to one contiguous
+  key/value range, LA Flash flattens the selected ranges, builds FlashAttention
+  `cu_seqlens_q` / `cu_seqlens_k`, and calls `flash_attn_varlen_func` directly.
+- **Packed MTP full-window plans:** for hybrid MTP decode, multiple full
+  key/value windows for the same query block are concatenated into one packed
+  key/value sequence before the FlashAttention call. This keeps the sparse
+  memory profile without constructing a `[B,H,Q,K]` attention mask.
+- **Compatible multi-segment plans:** when a query range attends to multiple
+  segments that cannot be packed as one sequence, each segment is evaluated with
+  FlashAttention and the partial outputs are merged with the standard
+  log-sum-exp softmax composition.
+
+The output tensor shape and dtype match the decoder attention output expected
+by the model. This path is inference-oriented and depends on FlashAttention's
+forward kernels; it is not a custom autograd training backend.
+
+## Runtime Knobs
+
+| Variable | Default | Meaning |
+| --- | --- | --- |
+| `LA_FLASH_ATTN` | `sdpa` | Set to `la_flash` to enable this backend through `batch_utils`. |
+| `LA_FLASH_FASTPATH` | `auto` | Use FlashAttention varlen for packed simple plans. |
+| `LA_FLASH_SEGMENT_FASTPATH` | `auto` | Use FlashAttention varlen for multi-segment sparse plans. Full segments are packed first; other compatible segments use LSE merging. |
+| `LA_FLASH_PLAN_STATS` | `0` | Record sparse plan statistics in inference summaries. |
+
+## Notes
+
+Dense prefill and stock worker-style generation should keep
+`LA_FLASH_DENSE_BACKEND=sdpa`; LA Flash is used for sparse range plans
+produced by `batch_utils`.
+
+This package is for inference and evaluation. Training remains on the
+MagiAttention backend; the batched sparse-plan decode runtime does not support
+the `labels` training path.
+
+## Source Layout
+
+- `range_attention.py`: FlashAttention varlen dispatch, sparse KV packing, LSE
+  merge fallback, and availability checks.

kernel_utils/__init__.py

@@ -0,0 +1,5 @@
+"""FlashAttention sparse range utilities for LocateAnything batch inference."""
+
+from .range_attention import range_attention, is_available
+
+__all__ = ["range_attention", "is_available"]

kernel_utils/range_attention.py

@@ -0,0 +1,394 @@
+"""Sparse LocateAnything attention implemented with FlashAttention varlen.
+
+The public API accepts flattened query/key/value tensors:
+
+  q: [total_q, num_q_heads, head_dim]
+  k: [total_k, num_kv_heads, head_dim]
+  v: [total_k, num_kv_heads, head_dim]
+
+and a Magi-style range plan:
+
+  q_ranges: [num_ranges, 2]
+  k_ranges: [num_key_segments, 2]
+  segment_offsets: [num_query_groups + 1]
+  attn_type_map:
+    0 = full attention over the listed key segment(s)
+    1 = bottom-right causal attention
+
+For LocateAnything hybrid MTP decode, batch_utils represents the window as a
+causal prefix plus full-attention sparse window segments.  This module packs
+those visible KV segments and calls FlashAttention varlen, avoiding dense masks.
+"""
+from __future__ import annotations
+
+import os
+from typing import Optional
+
+import torch
+
+
+_FLASH_ATTN_VARLEN = None
+_FLASH_ATTN_ERROR: Optional[BaseException] = None
+
+
+def _env_enabled(name: str, default: str = "auto") -> bool:
+    value = os.environ.get(name, default).strip().lower()
+    return value in {"", "auto", "1", "on", "true", "yes", "force"}
+
+
+def is_available() -> bool:
+    try:
+        _load_flash_attn_varlen()
+        return True
+    except Exception:
+        return False
+
+
+def _flash_fastpath_enabled() -> bool:
+    return _env_enabled("LA_FLASH_FASTPATH", "auto")
+
+
+def _flash_segment_fastpath_enabled() -> bool:
+    return _env_enabled("LA_FLASH_SEGMENT_FASTPATH", "auto")
+
+
+def _load_flash_attn_varlen():
+    global _FLASH_ATTN_VARLEN, _FLASH_ATTN_ERROR
+    if _FLASH_ATTN_VARLEN is not None:
+        return _FLASH_ATTN_VARLEN
+    if _FLASH_ATTN_ERROR is not None:
+        raise _FLASH_ATTN_ERROR
+    try:
+        from flash_attn import flash_attn_varlen_func
+
+        _FLASH_ATTN_VARLEN = flash_attn_varlen_func
+        return _FLASH_ATTN_VARLEN
+    except BaseException as exc:
+        _FLASH_ATTN_ERROR = exc
+        raise
+
+
+def _coalesce_query_groups(q_ranges, k_ranges, attn_type_map):
+    """Group consecutive entries that share the same query span and mask type."""
+    if q_ranges.numel() == 0:
+        segment_offsets = torch.zeros((1,), dtype=torch.int32, device=q_ranges.device)
+        return q_ranges, k_ranges, segment_offsets, attn_type_map, 0, 0
+
+    q_cpu = q_ranges.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    t_cpu = attn_type_map.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    grouped_q = []
+    grouped_t = []
+    offsets = [0]
+    max_q_len = 0
+    last_q = None
+    last_t = None
+    for idx, (qr, attn_type) in enumerate(zip(q_cpu.tolist(), t_cpu.tolist())):
+        key = (int(qr[0]), int(qr[1]))
+        attn_type = int(attn_type)
+        if attn_type not in (0, 1):
+            raise RuntimeError(
+                "LA Flash path only supports FlashAttention-compatible attn_type 0/1. "
+                f"Got attn_type={attn_type}; regenerate a type 0/1 range plan."
+            )
+        if last_q is None:
+            grouped_q.append([key[0], key[1]])
+            grouped_t.append(attn_type)
+            max_q_len = max(max_q_len, key[1] - key[0])
+            last_q = key
+            last_t = attn_type
+            continue
+        if key == last_q and attn_type == last_t:
+            continue
+        offsets.append(idx)
+        grouped_q.append([key[0], key[1]])
+        grouped_t.append(attn_type)
+        max_q_len = max(max_q_len, key[1] - key[0])
+        last_q = key
+        last_t = attn_type
+    offsets.append(int(q_ranges.shape[0]))
+
+    k_cpu = k_ranges.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    max_k_len = max((int(end) - int(start) for start, end in k_cpu.tolist()), default=0)
+
+    return (
+        torch.tensor(grouped_q, dtype=torch.int32, device=q_ranges.device).contiguous(),
+        k_ranges,
+        torch.tensor(offsets, dtype=torch.int32, device=q_ranges.device).contiguous(),
+        torch.tensor(grouped_t, dtype=torch.int32, device=q_ranges.device).contiguous(),
+        int(max_q_len),
+        int(max_k_len),
+    )
+
+
+def _flash_lse_to_tq_h(lse, total_q, q_lengths=None):
+    if lse is None:
+        return None
+    if lse.dim() != 2:
+        if lse.dim() == 3 and q_lengths is not None and lse.shape[0] == len(q_lengths):
+            chunks = []
+            for idx, q_len in enumerate(q_lengths):
+                q_len = int(q_len)
+                if lse.shape[1] == 0 or q_len > lse.shape[2]:
+                    return None
+                chunks.append(lse[idx, :, :q_len].transpose(0, 1).contiguous())
+            merged = torch.cat(chunks, dim=0).float()
+            return merged if merged.shape[0] == total_q else None
+        return None
+    if lse.shape[0] == total_q:
+        return lse.float()
+    if lse.shape[1] == total_q:
+        return lse.transpose(0, 1).contiguous().float()
+    return None
+
+
+def _make_cu_seqlens(lengths, device):
+    return torch.tensor([0] + list(torch.tensor(lengths).cumsum(0).tolist()), device=device, dtype=torch.int32)
+
+
+def _try_flash_segment_merge(
+    q,
+    k,
+    v,
+    k_ranges,
+    segment_offsets,
+    group_q_ranges,
+    group_attn_type_map,
+    softmax_scale,
+):
+    if not _flash_segment_fastpath_enabled():
+        return None
+    if q.dtype not in (torch.float16, torch.bfloat16) or k.dtype != q.dtype or v.dtype != q.dtype:
+        return None
+    if group_q_ranges is None or segment_offsets is None or group_attn_type_map is None:
+        return None
+
+    flash_attn_varlen = _load_flash_attn_varlen()
+    gq_cpu = group_q_ranges.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    kr_cpu = k_ranges.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    seg_cpu = segment_offsets.detach().to(device="cpu", dtype=torch.int32).contiguous()
+    type_cpu = group_attn_type_map.detach().to(device="cpu", dtype=torch.int32).contiguous()
+
+    groups = []
+    max_segments = 0
+    for group_idx, (q_start, q_end) in enumerate(gq_cpu.tolist()):
+        attn_type = int(type_cpu[group_idx].item())
+        if attn_type not in (0, 1):
+            return None
+        seg_start = int(seg_cpu[group_idx].item())
+        seg_end = int(seg_cpu[group_idx + 1].item())
+        if seg_end <= seg_start or q_end <= q_start:
+            return None
+        segments = kr_cpu[seg_start:seg_end].tolist()
+        max_segments = max(max_segments, len(segments))
+        groups.append((int(q_start), int(q_end), attn_type, [(int(a), int(b)) for a, b in segments]))
+
+    if not groups or max_segments == 0:
+        return None
+
+    can_pack_full_groups = all(attn_type == 0 or len(segments) == 1 for _, _, attn_type, segments in groups)
+    if can_pack_full_groups:
+        merged = torch.empty((q.shape[0], q.shape[1], q.shape[2]), device=q.device, dtype=q.dtype)
+        covered = torch.zeros((q.shape[0],), device=q.device, dtype=torch.bool)
+        for attn_type in (0, 1):
+            q_slices = []
+            k_slices = []
+            v_slices = []
+            q_lengths = []
+            k_lengths = []
+            targets = []
+            for q_start, q_end, group_type, segments in groups:
+                if group_type != attn_type:
+                    continue
+                q_slices.append(q[q_start:q_end])
+                if attn_type == 0 and len(segments) > 1:
+                    k_slices.append(torch.cat([k[start:end] for start, end in segments], dim=0))
+                    v_slices.append(torch.cat([v[start:end] for start, end in segments], dim=0))
+                    k_lengths.append(sum(end - start for start, end in segments))
+                else:
+                    k_start, k_end = segments[0]
+                    k_slices.append(k[k_start:k_end])
+                    v_slices.append(v[k_start:k_end])
+                    k_lengths.append(k_end - k_start)
+                q_lengths.append(q_end - q_start)
+                targets.append((q_start, q_end))
+            if not q_slices:
+                continue
+
+            out_pass = flash_attn_varlen(
+                torch.cat(q_slices, dim=0).contiguous(),
+                torch.cat(k_slices, dim=0).contiguous(),
+                torch.cat(v_slices, dim=0).contiguous(),
+                _make_cu_seqlens(q_lengths, q.device),
+                _make_cu_seqlens(k_lengths, q.device),
+                int(max(q_lengths)),
+                int(max(k_lengths)),
+                dropout_p=0.0,
+                softmax_scale=float(softmax_scale),
+                causal=bool(attn_type == 1),
+            )
+            if isinstance(out_pass, tuple):
+                out_pass = out_pass[0]
+
+            cursor = 0
+            for q_start, q_end in targets:
+                q_len = q_end - q_start
+                merged[q_start:q_end] = out_pass[cursor:cursor + q_len]
+                covered[q_start:q_end] = True
+                cursor += q_len
+
+        if bool(covered.all().item()):
+            return merged
+
+    merged = torch.zeros((q.shape[0], q.shape[1], q.shape[2]), device=q.device, dtype=torch.float32)
+    merged_lse = torch.full((q.shape[0], q.shape[1]), -float("inf"), device=q.device, dtype=torch.float32)
+    covered = torch.zeros((q.shape[0],), device=q.device, dtype=torch.bool)
+
+    for segment_idx in range(max_segments):
+        for attn_type in (0, 1):
+            q_slices = []
+            k_slices = []
+            v_slices = []
+            q_lengths = []
+            k_lengths = []
+            targets = []
+            for q_start, q_end, group_type, segments in groups:
+                if group_type != attn_type or segment_idx >= len(segments):
+                    continue
+                k_start, k_end = segments[segment_idx]
+                if k_end <= k_start:
+                    continue
+                q_slices.append(q[q_start:q_end])
+                k_slices.append(k[k_start:k_end])
+                v_slices.append(v[k_start:k_end])
+                q_lengths.append(q_end - q_start)
+                k_lengths.append(k_end - k_start)
+                targets.append((q_start, q_end))
+            if not q_slices:
+                continue
+
+            result = flash_attn_varlen(
+                torch.cat(q_slices, dim=0).contiguous(),
+                torch.cat(k_slices, dim=0).contiguous(),
+                torch.cat(v_slices, dim=0).contiguous(),
+                _make_cu_seqlens(q_lengths, q.device),
+                _make_cu_seqlens(k_lengths, q.device),
+                int(max(q_lengths)),
+                int(max(k_lengths)),
+                dropout_p=0.0,
+                softmax_scale=float(softmax_scale),
+                causal=bool(attn_type == 1),
+                return_attn_probs=True,
+            )
+            if not isinstance(result, tuple) or len(result) < 2:
+                return None
+            out_pass = result[0]
+            lse_pass = _flash_lse_to_tq_h(result[1], out_pass.shape[0], q_lengths)
+            if lse_pass is None:
+                return None
+
+            cursor = 0
+            for q_start, q_end in targets:
+                q_len = q_end - q_start
+                out_seg = out_pass[cursor:cursor + q_len].float()
+                lse_seg = lse_pass[cursor:cursor + q_len]
+                old_lse = merged_lse[q_start:q_end]
+                new_lse = torch.maximum(old_lse, lse_seg)
+                old_w = torch.exp(old_lse - new_lse)
+                seg_w = torch.exp(lse_seg - new_lse)
+                denom = (old_w + seg_w).clamp_min(1e-20)
+                merged[q_start:q_end] = (
+                    merged[q_start:q_end] * old_w.unsqueeze(-1)
+                    + out_seg * seg_w.unsqueeze(-1)
+                ) / denom.unsqueeze(-1)
+                merged_lse[q_start:q_end] = new_lse + torch.log(denom)
+                covered[q_start:q_end] = True
+                cursor += q_len
+
+    if not bool(covered.all().item()):
+        return None
+    return merged.to(dtype=q.dtype)
+
+
+def range_attention(
+    q,
+    k,
+    v,
+    q_ranges,
+    k_ranges,
+    attn_type_map,
+    softmax_scale: float,
+    *,
+    segment_offsets=None,
+    group_q_ranges=None,
+    group_attn_type_map=None,
+    max_q_len=None,
+    max_k_len=None,
+    flash_cu_seqlens_q=None,
+    flash_cu_seqlens_k=None,
+    flash_causal=None,
+    disjoint_q_ranges=None,
+):
+    """Run sparse range attention through FlashAttention varlen."""
+    del disjoint_q_ranges
+    if not q.is_cuda:
+        raise RuntimeError("LA Flash range_attention requires CUDA tensors")
+    if segment_offsets is None or group_q_ranges is None or group_attn_type_map is None:
+        (
+            group_q_ranges,
+            k_ranges,
+            segment_offsets,
+            group_attn_type_map,
+            computed_max_q_len,
+            computed_max_k_len,
+        ) = _coalesce_query_groups(q_ranges, k_ranges, attn_type_map)
+        if max_q_len is None:
+            max_q_len = computed_max_q_len
+        if max_k_len is None:
+            max_k_len = computed_max_k_len
+    elif max_q_len is None:
+        lengths = (group_q_ranges[:, 1] - group_q_ranges[:, 0]).detach().to(device="cpu")
+        max_q_len = int(lengths.max().item()) if lengths.numel() else 0
+    if max_k_len is None:
+        k_lengths = (k_ranges[:, 1] - k_ranges[:, 0]).detach().to(device="cpu")
+        max_k_len = int(k_lengths.max().item()) if k_lengths.numel() else 0
+
+    if (
+        flash_cu_seqlens_q is not None
+        and flash_cu_seqlens_k is not None
+        and flash_causal is not None
+        and _flash_fastpath_enabled()
+        and q.dtype in (torch.float16, torch.bfloat16)
+        and k.dtype == q.dtype
+        and v.dtype == q.dtype
+    ):
+        flash_attn_varlen = _load_flash_attn_varlen()
+        return flash_attn_varlen(
+            q.contiguous(),
+            k.contiguous(),
+            v.contiguous(),
+            flash_cu_seqlens_q.contiguous().to(device=q.device, dtype=torch.int32),
+            flash_cu_seqlens_k.contiguous().to(device=q.device, dtype=torch.int32),
+            int(max_q_len),
+            int(max_k_len),
+            dropout_p=0.0,
+            softmax_scale=float(softmax_scale),
+            causal=bool(flash_causal),
+        )
+
+    segment_out = _try_flash_segment_merge(
+        q,
+        k,
+        v,
+        k_ranges,
+        segment_offsets,
+        group_q_ranges,
+        group_attn_type_map,
+        softmax_scale,
+    )
+    if segment_out is not None:
+        return segment_out
+
+    raise RuntimeError(
+        "LA Flash could not express this range plan with FlashAttention varlen. "
+        "Only attn_type 0/1 range plans are supported in the release path."
+    )

mask_magi_utils.py

@@ -0,0 +1,101 @@
+# Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
+#
+# NVIDIA CORPORATION and its licensors retain all intellectual property
+# and proprietary rights in and to this software, related documentation
+# and any modifications thereto.  Any use, reproduction, disclosure or
+# distribution of this software and related documentation without an express
+# license agreement from NVIDIA CORPORATION is strictly prohibited.
+
+import torch
+
+# MagiAttention attn_type_map convention
+FULL, CAUSAL = 0, 1
+
+def build_magi_ranges(kv_len: int, q_len: int, block_size: int, ar_decode: bool=False, device: str = "cpu"):
+    """
+    Fixed strategy:
+      - use_cache=True: Mask blocked_k = (kv_len - block_size - 1) column
+      - causal_attn=False: Window interior is FULL (bidirectional)
+      - If q_len==kv_len: Use coarse prefix version (fewer ranges)
+      - Otherwise: General decode version (recompute rows expanding visible region row by row)
+
+    Conventions:
+      - K/V global length kv_len: [0, kv_len)
+      - Current Q is "last q_len tokens"
+      - First r=q_len-block_size rows are recomputed; last block_size rows are window
+    """
+    assert 0 < q_len <= kv_len
+
+    if ar_decode:
+        return {
+            "q_ranges": torch.tensor([[0, q_len]], dtype=torch.int32, device=device).contiguous(),
+            "k_ranges": torch.tensor([[0, kv_len]], dtype=torch.int32, device=device).contiguous(),
+            "attn_type_map": torch.tensor([CAUSAL],    dtype=torch.int32, device=device).contiguous(),
+        }
+
+
+    assert 0 < block_size <= q_len <= kv_len
+    B = block_size
+    r = q_len - B
+    q_global_start = kv_len - q_len
+
+    window_start_k = kv_len - B
+    blocked_k = window_start_k - 1  # The column that is blocked
+
+    q_ranges, k_ranges, types = [], [], []
+
+    # -------- prefix (q_len == kv_len) coarse-grained --------
+    if q_len == kv_len:
+        prefix_len = window_start_k  # kv_len - B
+
+        # prefix->prefix: causal
+        if prefix_len > 0:
+            q_ranges += [[0, prefix_len]]
+            k_ranges += [[0, prefix_len]]
+            types   += [CAUSAL]
+
+        # window->prefix: full, but exclude blocked_k => keys [0, blocked_k)
+        if prefix_len > 0 and blocked_k > 0:
+            q_ranges += [[prefix_len, kv_len]]
+            k_ranges += [[0, blocked_k]]
+            types   += [FULL]
+
+        # window->window: full
+        q_ranges += [[prefix_len, kv_len]]
+        k_ranges += [[prefix_len, kv_len]]
+        types += [FULL]
+
+        return {
+            "q_ranges": torch.tensor(q_ranges, dtype=torch.int32, device=device).contiguous(),
+            "k_ranges": torch.tensor(k_ranges, dtype=torch.int32, device=device).contiguous(),
+            "attn_type_map": torch.tensor(types,    dtype=torch.int32, device=device).contiguous(),
+        }
+
+    # -------- decode / general (q_len < kv_len) --------
+
+    # A) Recomputed rows: expand visible key cutoff row by row (use FULL + single-row q_range for precise shape)
+    for i in range(r):
+        g = q_global_start + i
+        q_ranges.append([i, i + 1])
+        k_ranges.append([0, g + 1])   # Allow keys [0, g]
+        types.append(FULL)
+
+    # B) Window rows: allow prefix but block blocked_k; window interior is full
+    q_win = [r, q_len]
+
+    # prefix keys [0, blocked_k)
+    if blocked_k > 0:
+        q_ranges.append(q_win)
+        k_ranges.append([0, blocked_k])
+        types.append(FULL)
+
+    # window keys [window_start_k, kv_len)
+    q_ranges.append(q_win)
+    k_ranges.append([window_start_k, kv_len])
+    types.append(FULL)
+
+    return {
+        "q_ranges": torch.tensor(q_ranges, dtype=torch.int32, device=device).contiguous(),
+        "k_ranges": torch.tensor(k_ranges, dtype=torch.int32, device=device).contiguous(),
+        "attn_type_map": torch.tensor(types, dtype=torch.int32, device=device).contiguous(),
+    }

mask_sdpa_utils.py

@@ -0,0 +1,232 @@
+# Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
+#
+# NVIDIA CORPORATION and its licensors retain all intellectual property
+# and proprietary rights in and to this software, related documentation
+# and any modifications thereto.  Any use, reproduction, disclosure or
+# distribution of this software and related documentation without an express
+# license agreement from NVIDIA CORPORATION is strictly prohibited.
+
+import torch
+
+
+def find_prefix_seq_length_by_pe(
+    pe: torch.Tensor
+) -> torch.Tensor:
+    """
+    Find the sequence length where position encoding drops (indicating prefix boundary).
+    Args:
+        pe: Position encoding tensor of shape [Batch size, Sequence length ]
+            Contains position indices for each token in the sequence.
+    Returns:
+        torch.Tensor: A tensor of shape [B] containing:
+            - The index where position encoding drops for each sequence
+            - -1 if no drop occurs in the sequence
+    """
+    batch_size, seq_len = pe.shape
+    prev = pe[:, :-1]
+    curr = pe[:, 1:]
+    drop_mask = curr < prev  #  [batch_size, seq_len-1]
+
+    seq_len = torch.full((batch_size,), -1, dtype=torch.long)
+    
+    for b in range(batch_size):
+        drop_pos = torch.nonzero(drop_mask[b], as_tuple=False)
+        if drop_pos.numel() > 0:
+            i = drop_pos[0].item() + 1  # Take first drop position (+1 because we compared shifted sequences)
+            seq_len[b] = i
+
+    return seq_len
+
+
+
+def update_causal_mask_with_pad_non_visible_2d(
+    input_ids: torch.Tensor,
+    attn_mask_2d: torch.Tensor,
+    text_mask_token_id: int,
+    block_size: int = 4,
+    causal_attn: bool = False
+) -> torch.Tensor:
+    """
+    Updates a 2D attention mask for hole sequence through input_ids and text_mask_token_id
+
+    Args:
+        input_ids: Input token IDs (unused in current implementation)
+        attn_mask_2d: 2D attention mask matrix of shape [seq_len, seq_len] where:
+            - 0.0 indicates allowed attention
+            - -inf indicates masked attention
+        text_mask_token_id: ID representing masked tokens
+        block_size: Size of the diffusion window
+        causal_attn: If True, maintains strict causal masking throughout
+    
+    Returns:
+        Modified attention mask with updated visibility patterns
+    """
+    seq_len = input_ids.shape[0]
+    device = input_ids.device
+
+    # Identify masked tokens and their preceding positions
+    input_mask = input_ids.eq(text_mask_token_id)
+    input_before_mask = torch.zeros_like(input_mask)
+    input_before_mask[:-1] = input_mask[1:]
+    mask_cols = (input_mask | input_before_mask)
+    non_mask = ~mask_cols
+
+    rows = torch.arange(seq_len, device=device)[:, None]
+    cols = torch.arange(seq_len, device=device)
+
+    indices = torch.arange(seq_len, device=device)
+    prev_non_mask = (indices * non_mask).cummax(dim=0).values
+
+    max_value = torch.iinfo(indices.dtype).max
+    mask_indices = torch.where(non_mask, indices, torch.full_like(indices, max_value))
+    reversed_mask_indices = torch.flip(mask_indices, dims=[0])
+    reversed_cummin = reversed_mask_indices.cummin(dim=0).values
+    next_non_mask = torch.flip(reversed_cummin, dims=[0])
+
+    infra_mask = (
+        (cols > prev_non_mask) &
+        (rows >= next_non_mask[None, :]) &
+        mask_cols[None, :]
+    )
+    attn_mask_2d.masked_fill_(infra_mask, -float('inf'))
+
+    if not causal_attn:
+        visible_mask = (
+            (rows > prev_non_mask[None, :]) &
+            (rows < cols) &
+            mask_cols[None, :]
+        )
+        attn_mask_2d.masked_fill_(visible_mask, 0.0)
+
+    return attn_mask_2d
+
+
+def update_causal_mask_for_one_gen_window_2d(
+    input_ids: torch.Tensor,
+    attn_mask_2d: torch.Tensor,
+    block_size: int = 4,
+    use_cache: bool = True,
+    causal_attn: bool = False
+) -> torch.Tensor:
+    """
+    Updates a 2D attention mask for a diffusion window in transformer inference.
+
+    Args:
+        input_ids: Input token IDs (unused in current implementation)
+        attn_mask_2d: 2D attention mask matrix of shape [seq_len, seq_len] where:
+            - 0.0 indicates allowed attention
+            - -inf indicates masked attention
+        block_size: Size of the diffusion window
+        use_cache: Whether key-value cache is being used
+        causal_attn: If True, maintains strict causal masking throughout
+    
+    Returns:
+        Modified attention mask with updated visibility patterns
+    """
+
+    if not causal_attn:
+        # Make the diffusion window (last block_size tokens) fully visible to itself
+        # This allows bidirectional attention within the diffusion window
+        attn_mask_2d[-block_size:, -block_size:] = 0.0
+    if use_cache:
+        # Mask the last token from previous round to prevent recomputation and maintain generation consistency.
+        attn_mask_2d[-block_size:, -block_size-1] = -float('inf')
+
+    return attn_mask_2d
+
+
+def create_block_diff_mask_by_pe_4d(
+    block_size: int, 
+    x0_len_list: torch.Tensor, 
+    position_ids: torch.Tensor, 
+    causal_attn: bool = False
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """Generates a 4D attention mask for block-difference attention patterns.
+
+    The mask consists of three regions:
+    1. Causal block (top-left): Standard causal attention for `x0` tokens.
+    2. Mutual block (bottom-right): Non-causal attention within the same block for non-`x0` tokens.
+    3. Prefix block (bottom-left): Non-`x0` tokens can attend to a prefix of `x0` tokens.
+
+    Args:
+        block_size (int): Size of processing blocks for non-`x0` tokens.
+        x0_len_list (torch.Tensor): Tensor of shape [B] containing lengths of `x0` segments per batch.
+        position_ids (torch.Tensor): Tensor of shape [B, seq_len] containing position IDs.
+        causal_attn (bool, optional): If True, enforces causal masking in mutual blocks. Defaults to False.
+
+    Returns:
+        tuple[torch.Tensor, torch.Tensor]:
+            - A float mask of shape [batch_size, 1, seq_len, seq_len] with `-inf` for masked positions (non visiable).
+            - A boolean mask of shape [batch_size, 1, seq_len, seq_len] indicating allowed attention positions.
+    """
+    batch_size, seq_len = position_ids.shape
+    device = position_ids.device
+    
+    # Create position indices [batch_size, seq_len, seq_len]
+    q_idx = torch.arange(seq_len, device=device).view(1, seq_len, 1)  # [1, seq_len, 1]
+    kv_idx = torch.arange(seq_len, device=device).view(1, 1, seq_len)  # [1, 1, seq_len]
+    
+    # Broadcast to [B, seq_len, seq_len]
+    x0_len = x0_len_list.view(batch_size, 1, 1)  # [batch_size, 1, 1]
+    x0_flag_q = q_idx < x0_len           # [batch_size, seq_len, seq_len]
+    x0_flag_kv = kv_idx < x0_len         
+    
+    # Block indices calculation [batch_size, seq_len, seq_len]
+    q_block_idx = (q_idx - x0_len) // block_size
+    kv_block_idx = (kv_idx - x0_len) // block_size
+    
+    # causal block (top-left)
+    block_causal = x0_flag_q & x0_flag_kv & (q_idx >= kv_idx)
+
+    mutual_condition = (q_idx >= kv_idx) if causal_attn else torch.ones_like(q_idx, dtype=torch.bool)
+    block_mutual = (
+        ~x0_flag_q & ~x0_flag_kv &
+        (q_block_idx == kv_block_idx) &
+        mutual_condition
+    )
+
+    q_blk = torch.div(q_idx - x0_len, block_size, rounding_mode='floor')
+    q_blk_start = (x0_len_list.view(batch_size, 1) + q_blk[:, :, 0] * block_size).clamp(min=0, max=seq_len - 1)
+    prefix_len = position_ids.gather(1, q_blk_start)
+    prefix_len = prefix_len.unsqueeze(2)
+    block_prefix = (~x0_flag_q & x0_flag_kv) & (kv_idx < prefix_len)
+
+    final_mask = (block_causal | block_mutual | block_prefix)
+    customized_mask = torch.full_like(final_mask, float('-inf'), dtype=torch.bfloat16)
+    customized_mask.masked_fill_(final_mask, 0.0)
+    
+    return customized_mask.unsqueeze(1).to(device=device), final_mask.unsqueeze(1).to(device=device)
+
+
+def find_pred_pos_from_input_ids(
+    input_ids: torch.LongTensor = None,
+    text_mask_token_id: int = None,
+) -> torch.Tensor:
+    """Compute the relative prediction positions for masked tokens in a sequence.
+
+    For non-masked positions, the output is 0. For masked positions, the value increments
+    by 1 for each consecutive mask token, indicating how many steps ahead the prediction is.
+
+    Args:
+        input_ids (torch.LongTensor): Input token IDs of shape [batch_size, seq_len].
+        text_mask_token_id (int, optional): Token ID representing masked positions. Defaults to 151666.
+
+    Returns:
+        torch.Tensor: A tensor of shape [batch_size, seq_len] where:
+            - 0 indicates a non-masked token.
+            - n > 0 indicates the nth consecutive masked token (e.g., 1 = first mask, 2 = second mask, etc.).
+    """
+    batch_size, seq_len = input_ids.shape
+    device = input_ids.device
+
+    is_mask = (input_ids == text_mask_token_id)
+
+    base_mask = torch.zeros((batch_size, seq_len), dtype=torch.int8, device=device)
+
+    for b in range(batch_size):
+        for ix in range(1, seq_len):
+            if is_mask[b][ix] == True:
+                # Increment counter if current token is masked
+                base_mask[b][ix] = base_mask[b][ix-1] + 1
+
+    return base_mask

model-00001-of-00002.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:923cfc10fed19808067da6df85a9a4220ddc1f9eb91ceee94c0fecd05d0f2d58
+size 4959632160

model-00002-of-00002.safetensors

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:3459ba101f40594f3f62d3312014f1f8378b4ba3da3b1d562480045938fc7d47
+size 2701795216

model.safetensors.index.json

@@ -0,0 +1,777 @@
+{
+  "metadata": {
+    "total_size": 7661331936
+  },
+  "weight_map": {
+    "language_model.lm_head.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.embed_tokens.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.0.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.1.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.10.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.11.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.12.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.13.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.14.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.15.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.16.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.17.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.18.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.19.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.2.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.20.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.21.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.22.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.22.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.22.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.23.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.23.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.24.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.25.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.26.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.27.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.28.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.29.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.3.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.3.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.30.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.30.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.31.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.32.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.33.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.34.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.input_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.mlp.down_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.mlp.gate_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.mlp.up_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.post_attention_layernorm.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.k_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.k_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.o_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.q_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.q_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.v_proj.bias": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.35.self_attn.v_proj.weight": "model-00002-of-00002.safetensors",
+    "language_model.model.layers.4.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.4.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.5.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.6.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.7.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.8.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.input_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.mlp.down_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.mlp.gate_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.mlp.up_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.post_attention_layernorm.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.k_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.k_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.o_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.q_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.q_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.v_proj.bias": "model-00001-of-00002.safetensors",
+    "language_model.model.layers.9.self_attn.v_proj.weight": "model-00001-of-00002.safetensors",
+    "language_model.model.norm.weight": "model-00002-of-00002.safetensors",
+    "mlp1.0.bias": "model-00002-of-00002.safetensors",
+    "mlp1.0.weight": "model-00002-of-00002.safetensors",
+    "mlp1.1.bias": "model-00002-of-00002.safetensors",
+    "mlp1.1.weight": "model-00002-of-00002.safetensors",
+    "mlp1.3.bias": "model-00002-of-00002.safetensors",
+    "mlp1.3.weight": "model-00002-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.0.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.1.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.10.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.11.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.12.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.13.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.14.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.15.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.16.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.17.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.18.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.19.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.2.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.20.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.21.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.22.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.23.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.24.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.25.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.26.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.3.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.4.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.5.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.6.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.7.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.8.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.mlp.fc0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.mlp.fc0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.mlp.fc1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.mlp.fc1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.norm0.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.norm0.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.norm1.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.norm1.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.wo.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.wo.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.wqkv.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.blocks.9.wqkv.weight": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.final_layernorm.bias": "model-00001-of-00002.safetensors",
+    "vision_model.encoder.final_layernorm.weight": "model-00001-of-00002.safetensors",
+    "vision_model.patch_embed.pos_emb.weight": "model-00001-of-00002.safetensors",
+    "vision_model.patch_embed.proj.bias": "model-00001-of-00002.safetensors",
+    "vision_model.patch_embed.proj.weight": "model-00001-of-00002.safetensors"
+  }
+}

modeling_locateanything.py

@@ -0,0 +1,537 @@
+# --------------------------------------------------------
+# NVIDIA
+# Copyright (c) 2025 NVIDIA
+# Licensed under The MIT License [see LICENSE for details]
+# --------------------------------------------------------
+
+import time
+from typing import List, Optional, Tuple, Union
+
+import numpy as np
+import torch
+from torch import nn
+from torch.nn import CrossEntropyLoss
+from transformers.generation import GenerationMixin
+from transformers.modeling_outputs import CausalLMOutputWithPast
+from transformers.modeling_utils import PreTrainedModel
+from transformers.utils import add_start_docstrings, is_flash_attn_2_available, logging
+from peft import LoraConfig, get_peft_model
+
+from .configuration_locateanything import LocateAnythingConfig
+from .modeling_qwen2 import Qwen2ForCausalLM
+from .modeling_vit import MoonVitPretrainedModel
+from transformers.models.qwen3.modeling_qwen3 import Qwen3ForCausalLM
+from .mask_sdpa_utils import *
+from .mask_magi_utils import *
+from .configuration_qwen2 import Qwen2Config
+
+from .generate_utils import (
+    sample_tokens,
+    handle_pattern,
+    get_token_ids_from_config,
+)
+
+logger = logging.get_logger(__name__)
+
+
+LOCATEANYTHING_START_DOCSTRING = r"""
+    This model inherits from [`PreTrainedModel`]. Check the superclass documentation for the generic methods the
+    library implements for all its model (such as downloading or saving, resizing the input embeddings, pruning heads
+    etc.)
+    This model is also a PyTorch [torch.nn.Module](https://pytorch.org/docs/stable/nn.html#torch.nn.Module) subclass.
+    Use it as a regular PyTorch Module and refer to the PyTorch documentation for all matter related to general usage
+    and behavior.
+    Parameters:
+        config ([`LocateAnythingConfig`]):
+            Model configuration class with all the parameters of the model. Initializing with a config file does not
+            load the weights associated with the model, only the configuration. Check out the
+            [`~PreTrainedModel.from_pretrained`] method to load the model weights.
+"""
+
+@add_start_docstrings(
+    "The bare LocateAnything Model outputting raw hidden-states without any specific head on top.",
+    LOCATEANYTHING_START_DOCSTRING,
+)
+class LocateAnythingPreTrainedModel(PreTrainedModel):
+    config_class = LocateAnythingConfig
+    base_model_prefix = "model"
+    main_input_name = 'input_ids'
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["Qwen2DecoderLayer"]
+    _skip_keys_device_placement = "past_key_values"
+    _supports_flash_attn_2 = True
+    _supports_cache_class = True
+    _supports_static_cache = True
+    _supports_quantized_cache = True
+    _supports_sdpa = True
+
+    @classmethod
+    def _autoset_attn_implementation(cls, config, *args, **kwargs):
+        if getattr(config, '_attn_implementation', None) == 'magi':
+            return config
+        return super()._autoset_attn_implementation(config, *args, **kwargs)
+
+    def _check_and_adjust_attn_implementation(self, attn_implementation, is_init_check=False):
+        if attn_implementation == "magi":
+            return "magi"
+        return super()._check_and_adjust_attn_implementation(attn_implementation, is_init_check)
+    
+    def _init_weights(self, module):
+        std = getattr(self.config, 'initializer_range', None) or self.config.text_config.initializer_range
+        if isinstance(module, (nn.Linear, nn.Conv2d)):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.bias is not None:
+                module.bias.data.zero_()
+        elif isinstance(module, nn.Embedding):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.padding_idx is not None:
+                module.weight.data[module.padding_idx].zero_()
+
+
+class LocateAnythingForConditionalGeneration(LocateAnythingPreTrainedModel, GenerationMixin):
+    config_class = LocateAnythingConfig
+    def __init__(self, config: LocateAnythingConfig, vision_model=None, language_model=None):
+        super().__init__(config)
+
+        self.template = config.template
+        self.mlp_checkpoint = config.mlp_checkpoint
+
+        logger.info(f'mlp_checkpoint: {self.mlp_checkpoint}')
+        if vision_model is not None:
+            self.vision_model = vision_model
+        else:
+            if config.vision_config.model_type == 'moonvit':
+                vision_attn_impl = getattr(config.vision_config, '_attn_implementation', None) or 'flash_attention_2'
+                if vision_attn_impl == 'flash_attention_2' and not is_flash_attn_2_available():
+                    logger.warning_once(
+                        "flash_attn is not available for MoonViT inference; falling back to sdpa."
+                    )
+                    vision_attn_impl = 'sdpa'
+                config.vision_config._attn_implementation = vision_attn_impl
+                self.vision_model = MoonVitPretrainedModel(config.vision_config)
+            else:
+                raise ValueError(f'Unsupported vision model type: {config.vision_config.model_type}. Only moonvit is supported.')
+
+        text_attn_impl = (
+            getattr(config.text_config, '_attn_implementation', None)
+            or getattr(config, '_attn_implementation', None)
+            or 'magi'
+        )
+        config.text_config._attn_implementation = text_attn_impl
+
+        if language_model is not None:
+            self.language_model = language_model
+        else:
+            if config.text_config.architectures[0] == 'Qwen2ForCausalLM':
+                self.language_model = Qwen2ForCausalLM(config.text_config)
+            elif config.text_config.architectures[0] == 'Qwen3ForCausalLM':
+                self.language_model = Qwen3ForCausalLM(config.text_config)
+            else:
+                raise ValueError(f'Unsupported language model architecture: {config.text_config.architectures[0]}. Only Qwen2ForCausalLM and Qwen3ForCausalLM are supported.')
+
+        vit_hidden_size = config.vision_config.hidden_size
+        llm_hidden_size = config.text_config.hidden_size
+
+        # MLP for moonvit (without pixel_shuffle_back, direct mapping)
+        self.mlp1 = nn.Sequential(
+                nn.LayerNorm(vit_hidden_size*4),
+                nn.Linear(vit_hidden_size*4, llm_hidden_size),
+                nn.GELU(),
+                nn.Linear(llm_hidden_size, llm_hidden_size)
+            )
+        self.image_token_index = config.image_token_index
+        self.neftune_alpha = None
+
+        if config.use_backbone_lora:
+            self.wrap_backbone_lora(r=config.use_backbone_lora, lora_alpha=2 * config.use_backbone_lora)
+
+        self.use_llm_lora = config.use_llm_lora
+        if config.use_llm_lora:
+            self.wrap_llm_lora(r=config.use_llm_lora, lora_alpha=2 * config.use_llm_lora)
+
+        self.token_ids = get_token_ids_from_config(config)
+
+        # Set _no_split_modules dynamically based on the actual LLM architecture
+        arch = config.text_config.architectures[0] if hasattr(config.text_config, 'architectures') and config.text_config.architectures else 'Qwen2ForCausalLM'
+        if 'Qwen3' in arch:
+            self._no_split_modules = ["Qwen3DecoderLayer"]
+        else:
+            self._no_split_modules = ["Qwen2DecoderLayer"]
+
+        
+    def wrap_backbone_lora(self, r=128, lora_alpha=256, lora_dropout=0.05):
+        lora_config = LoraConfig(
+            r=r,
+            target_modules=['self_attn.q_proj', 'self_attn.k_proj', 'self_attn.v_proj', 'self_attn.out_proj',
+                            'mlp.fc1', 'mlp.fc2'],
+            lora_alpha=lora_alpha,
+            lora_dropout=lora_dropout,
+        )
+        self.vision_model = get_peft_model(self.vision_model, lora_config)
+        self.vision_model.print_trainable_parameters()
+
+    def wrap_llm_lora(self, r=128, lora_alpha=256, lora_dropout=0.05):
+        lora_config = LoraConfig(
+            r=r,
+            target_modules=['self_attn.q_proj', 'self_attn.k_proj', 'self_attn.v_proj', 'self_attn.o_proj',
+                            'mlp.gate_proj', 'mlp.down_proj', 'mlp.up_proj'],
+            lora_alpha=lora_alpha,
+            lora_dropout=lora_dropout,
+            task_type='CAUSAL_LM'
+        )
+        self.language_model = get_peft_model(self.language_model, lora_config)
+        self.language_model.enable_input_require_grads()
+        self.language_model.print_trainable_parameters()
+        self.use_llm_lora = True
+        
+    
+    def forward(
+            self,
+            pixel_values: List[torch.FloatTensor],
+            input_ids: torch.LongTensor = None,
+            attention_mask: Optional[torch.Tensor] = None,
+            position_ids: Optional[torch.LongTensor] = None,
+            image_grid_hws: Optional[torch.Tensor] = None,
+            image_flags: Optional[torch.Tensor] = None,
+            past_key_values: Optional[List[torch.FloatTensor]] = None,
+            labels: Optional[torch.LongTensor] = None,
+            use_cache: Optional[bool] = None,
+            output_attentions: Optional[bool] = None,
+            output_hidden_states: Optional[bool] = None,
+            return_dict: Optional[bool] = None,
+    ) -> Union[Tuple, CausalLMOutputWithPast]:
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        input_embeds = self.language_model.get_input_embeddings()(input_ids)
+
+        has_images = image_flags is not None and image_flags.sum() > 0
+        
+        vit_embeds = self.extract_feature(pixel_values, image_grid_hws)
+            
+        B, N, C = input_embeds.shape
+        input_embeds = input_embeds.reshape(B * N, C)
+
+        if has_images:
+            filtered_vit_embeds = []
+            idx = 0
+            for flag in image_flags:
+                flag_val = flag.item()
+                if flag_val != 0:
+                    filtered_vit_embeds.extend(vit_embeds[idx:idx + flag_val])
+                    idx += flag_val
+                else:
+                    idx += 1
+
+            vit_embeds = filtered_vit_embeds
+            vit_embeds = torch.cat(vit_embeds, dim=0)
+
+            vit_embeds = self.mlp1(vit_embeds)
+            input_ids = input_ids.reshape(B * N)
+            selected = (input_ids == self.image_token_index)
+
+            input_embeds[selected] = input_embeds[selected] * 0.0 + vit_embeds[:selected.sum()]
+        else:
+            if vit_embeds:
+                vit_embeds = torch.cat(vit_embeds, dim=0)
+                vit_embeds = self.mlp1(vit_embeds)
+                input_ids = input_ids.reshape(B * N)
+                selected = (input_ids == self.image_token_index)
+                if selected.sum() > 0:
+                    input_embeds[selected] = vit_embeds[:selected.sum()]
+
+        input_embeds = input_embeds.reshape(B, N, C)
+
+        outputs = self.language_model(
+            inputs_embeds=input_embeds,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_values=past_key_values,
+            use_cache=use_cache,
+            output_attentions=output_attentions,
+            output_hidden_states=output_hidden_states,
+        )
+        logits = outputs.logits
+
+        loss = None
+        if labels is not None:
+            # Shift so that tokens < n predict n
+            shift_logits = logits[..., :-1, :].contiguous()
+            shift_labels = labels[..., 1:].contiguous()
+            # Flatten the tokens
+            loss_fct = CrossEntropyLoss()
+            shift_logits = shift_logits.view(-1, self.language_model.config.vocab_size)
+            shift_labels = shift_labels.view(-1)
+            # Enable model parallelism
+            shift_labels = shift_labels.to(shift_logits.device)
+            loss = loss_fct(shift_logits, shift_labels)
+
+        if not return_dict:
+            output = (logits,) + outputs[1:]
+            return (loss,) + output if loss is not None else output
+
+        return CausalLMOutputWithPast(
+            loss=loss,
+            logits=logits,
+            past_key_values=outputs.past_key_values,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )
+
+    
+    def extract_feature(self, pixel_values, image_grid_hws):
+        vit_embeds = self.vision_model(pixel_values=pixel_values, grid_hws=image_grid_hws)
+
+        return vit_embeds
+
+    def get_input_embeddings(self):
+        return self.language_model.get_input_embeddings()
+
+    def set_input_embeddings(self, value):
+        self.language_model.set_input_embeddings(value)
+
+    def get_output_embeddings(self):
+        return self.language_model.get_output_embeddings()
+
+    def set_output_embeddings(self, new_embeddings):
+        self.language_model.set_output_embeddings(new_embeddings)
+
+    def set_decoder(self, decoder):
+        self.language_model.set_decoder(decoder)
+
+    def get_decoder(self):
+        return self.language_model.get_decoder()
+
+    @torch.no_grad()
+    def generate(
+        self,
+        pixel_values: Optional[torch.FloatTensor] = None,
+        input_ids: Optional[torch.FloatTensor] = None,
+        attention_mask: Optional[torch.LongTensor] = None,
+        visual_features: Optional[torch.FloatTensor] = None,
+        image_grid_hws: Optional[torch.Tensor] = None,
+        tokenizer = None, 
+        n_future_tokens: int = 6,
+        **generate_kwargs,
+    ) -> torch.LongTensor:
+
+        verbose = generate_kwargs.pop('verbose', False)
+        start_time = time.time()
+        prefill_time = None
+
+        pixel_values = pixel_values.to(self.language_model.dtype)
+        # Convert numpy array to tensor if needed
+        if isinstance(image_grid_hws, np.ndarray):
+            image_grid_hws = torch.from_numpy(image_grid_hws).to(pixel_values.device, dtype=torch.int32)
+
+        batch_size, seq_len = input_ids.shape
+        assert batch_size == 1, 'only batch size = 1 is supported now'
+        assert generate_kwargs.get('use_cache', False), "Only use_cache=True is supported."
+
+        generated = input_ids.clone()
+        total_gen_length = min(tokenizer.model_max_length, seq_len + generate_kwargs.get('max_new_tokens', 2048))
+        iter_round = 0
+        past_key_values = None
+
+        # Extract visual features once before the loop
+        if visual_features is not None:
+            vit_embeds = visual_features
+        elif pixel_values is not None:
+            vit_embeds = self.extract_feature(pixel_values, image_grid_hws)
+        else:
+            vit_embeds = None
+        
+        if image_grid_hws is not None:
+            vit_embeds = torch.cat(vit_embeds, dim=0)
+            vit_embeds = self.mlp1(vit_embeds)
+
+        # ==================== Generation Mode ====================
+        # 'fast'   : MTP only, never fall back to AR
+        # 'slow'   : AR only, pure auto-regressive decoding
+        # 'hybrid' : MTP first, fall back to AR on error, switch back on box_end
+        generation_mode = generate_kwargs.get('generation_mode', 'hybrid')
+        assert generation_mode in ('fast', 'slow', 'hybrid'), \
+            f"Unsupported generation_mode='{generation_mode}'. Use 'fast', 'slow', or 'hybrid'."
+
+        sampling_history = []
+
+
+        use_mtp = generation_mode in ('fast', 'hybrid')
+        switch_to_ar_count = 0
+
+        # Pre-allocate mask tokens and position ids
+        default_mask_token_id = self.token_ids['default_mask_token_id']
+        pre_mask_tokens = torch.full(
+            (batch_size, n_future_tokens - 1),
+            default_mask_token_id,
+            dtype=generated.dtype,
+            device=generated.device
+        )
+        max_possible_len = total_gen_length + n_future_tokens
+        full_position_ids = torch.arange(0, max_possible_len, device=generated.device).unsqueeze(0)
+
+
+        def _prepare_inputs_in_mtp(generated):
+            generated_with_mask = torch.cat(
+                (
+                    generated, 
+                    generated[:, -1].unsqueeze(1),
+                    pre_mask_tokens
+                ), 
+                dim=1
+            ) # [batch_size, seq_len + 1 +  n_future_tokens - 1]
+
+            # Update pe for kvcache
+            start_idx = past_key_values[0][0].size(2) if past_key_values is not None else 0
+            position_ids = full_position_ids[:, start_idx : generated_with_mask.size(1)].clone()
+            position_ids[0, -n_future_tokens:] -= 1
+
+            prepare_inputs = self.language_model.prepare_inputs_for_generation(
+                generated_with_mask,
+                past_key_values,
+                None,
+                inputs_embeds=None,
+                use_cache=True,
+                position_ids=position_ids
+            )
+            return prepare_inputs
+
+
+        def _prepare_input_in_ar(generated):
+            start_idx = past_key_values[0][0].size(2) if past_key_values is not None else 0
+            position_ids = full_position_ids[:, start_idx : generated.size(1)]
+            prepare_inputs = self.language_model.prepare_inputs_for_generation(
+                generated,
+                past_key_values,
+                None,
+                inputs_embeds=None,
+                use_cache=True,
+                position_ids=position_ids
+            )
+            return prepare_inputs
+
+
+        def _sample_token_in_mtp(generated, outputs):
+            """Sample tokens using MTP (Multi-Token Prediction) mode."""
+            next_token_logits = outputs.logits[:, -n_future_tokens:, :]
+            probs, confidence, x0, box_avg = sample_tokens(
+                next_token_logits, generated, self.token_ids, keep_k=5, **generate_kwargs
+            )
+
+            is_box_empty = (box_avg[0] == 0).all()
+            new_tokens = x0[0] if is_box_empty else box_avg[0]
+
+            out_pattern = handle_pattern(new_tokens, self.token_ids, generation_mode)
+            out_type = out_pattern['type']
+            out_token = torch.tensor(out_pattern['tokens'], dtype=x0.dtype, device=x0.device)
+
+            return out_type, out_token
+
+
+        def _sample_token_in_ar(generated, outputs):
+            """Sample a single token using AR (Auto-Regressive) mode."""
+            next_token_logits = outputs.logits[:, -1:, :]
+            probs, confidence, x0, _ = sample_tokens(
+                next_token_logits, generated, self.token_ids, **generate_kwargs
+            )
+
+            out_token = x0[0]
+            out_type = 'continue_ar'
+            token_val = out_token[0].item()
+
+            box_end_token_id = self.token_ids['box_end_token_id']
+            coord_start_token_id = self.token_ids['coord_start_token_id']
+            coord_end_token_id = self.token_ids['coord_end_token_id']
+            none_token_id = self.token_ids['none_token_id']
+            im_end_token_id = self.token_ids['im_end_token_id']
+
+            if generation_mode == 'hybrid':
+                # Hybrid AR phase: detect box boundaries to switch back to MTP
+                if token_val == box_end_token_id:
+                    out_type = 'box_end_ar'
+                elif coord_start_token_id <= token_val <= coord_end_token_id or token_val == none_token_id:
+                    out_type = 'coord_ar'
+                else:
+                    out_type = 'im_end'
+            else:
+                # Slow mode: pure AR, only stop on im_end
+                if token_val == im_end_token_id:
+                    out_type = 'im_end'
+
+            return out_type, out_token
+
+
+        # Generate loop
+        while generated.size(1) < total_gen_length:
+            iter_round += 1
+
+            # Step 1: Prepare inputs
+            if use_mtp:
+                prepare_inputs = _prepare_inputs_in_mtp(generated)
+            else:
+                prepare_inputs = _prepare_input_in_ar(generated)
+
+            if iter_round == 1:
+                prepare_inputs.update({
+                    'visual_features': vit_embeds,
+                    'image_token_index': self.config.image_token_index,
+                })
+
+            # Step 2: Model forward & update KV cache
+            with torch.no_grad():
+                outputs = self.language_model(**prepare_inputs)
+
+            past_key_values = tuple(
+                (kv[0][:, :, :generated.shape[1], :], kv[1][:, :, :generated.shape[1], :])
+                for kv in outputs.past_key_values
+            )
+
+            # Step 3: Sample tokens
+            if use_mtp:
+                out_type, out_token = _sample_token_in_mtp(generated, outputs)
+            else:
+                out_type, out_token = _sample_token_in_ar(generated, outputs)
+
+            if verbose:
+                sampling_history.append(('ar' if 'ar' in out_type else 'mtp', tokenizer.decode(out_token, skip_special_tokens=False)))
+
+            generated = torch.cat([generated, out_token.unsqueeze(0)], dim=1)
+
+            # Step 4: Mode switching & termination
+            if out_type == 'im_end':
+                break
+
+            if generation_mode == 'hybrid':
+                if out_type == 'error_box':
+                    use_mtp = False
+                    switch_to_ar_count += 1
+                elif out_type == 'box_end_ar':
+                    use_mtp = True
+            # fast mode: use_mtp stays True always
+            # slow mode: use_mtp stays False always
+
+            if prefill_time is None:
+                prefill_time = time.time() - start_time
+
+        # Decode and return
+        generated_ids = generated[:, seq_len:]
+        response = tokenizer.batch_decode(generated_ids, skip_special_tokens=False)
+
+        if verbose:
+            end_time = time.time()
+            num_tokens = generated_ids.size(1)
+            num_boxes = response[0].count("<box>")
+            total_time = end_time - start_time
+
+            out_info =  f"\nStatistic Info, num_tokens={num_tokens}; " + \
+                    f"generate_time(s)={total_time:.4f}; " + \
+                    f"tps={(num_tokens / total_time):.4f}; " + \
+                    f"forward_step={iter_round}; " + \
+                    f"num_boxes={num_boxes}; " + \
+                    f"bps={(num_boxes / total_time):.4f}; " + \
+                    f"prefill_time={(prefill_time):.4f}; " + \
+                    f"switch_to_ar={switch_to_ar_count}\n"
+            print(out_info)
+
+            return response[0], sampling_history, out_info
+
+        return response[0]

modeling_qwen2.py

@@ -0,0 +1,1738 @@
+# coding=utf-8
+# Copyright 2024 The Qwen team, Alibaba Group and the HuggingFace Inc. team. All rights reserved.
+#
+# This code is based on EleutherAI's GPT-NeoX library and the GPT-NeoX
+# and OPT implementations in this library. It has been modified from its
+# original forms to accommodate minor architectural differences compared
+# to GPT-NeoX and OPT used by the Meta AI team that trained the model.
+#
+# 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 Qwen2 model."""
+import inspect
+import math
+import copy
+import warnings
+from functools import partial
+from typing import List, Optional, Tuple, Union
+
+import torch
+import torch.nn.functional as F
+import torch.utils.checkpoint
+from torch import nn
+from torch.nn import BCEWithLogitsLoss, CrossEntropyLoss, MSELoss
+
+from transformers.activations import ACT2FN
+from transformers.cache_utils import Cache, DynamicCache
+from transformers.modeling_attn_mask_utils import _prepare_4d_causal_attention_mask, _prepare_4d_causal_attention_mask_for_sdpa
+from transformers.modeling_outputs import BaseModelOutputWithPast, CausalLMOutputWithPast, SequenceClassifierOutputWithPast
+from transformers.modeling_utils import PreTrainedModel
+from transformers.utils import (
+    add_start_docstrings,
+    add_start_docstrings_to_model_forward,
+    is_flash_attn_2_available,
+    is_flash_attn_greater_or_equal_2_10,
+    logging,
+    replace_return_docstrings,
+)
+from .configuration_qwen2 import Qwen2Config
+
+if is_flash_attn_2_available():
+    from flash_attn import flash_attn_func, flash_attn_varlen_func
+    from flash_attn.bert_padding import index_first_axis, pad_input, unpad_input  # noqa
+
+    _flash_supports_window_size = "window_size" in list(inspect.signature(flash_attn_func).parameters)
+
+
+logger = logging.get_logger(__name__)
+
+# Magi Attention Supported
+_MAGI_AVAILABLE = False
+try:
+    from magi_attention.functional.flex_flash_attn import flex_flash_attn_func
+    _MAGI_AVAILABLE = True
+except ImportError:
+    flex_flash_attn_func = None
+
+
+_CHECKPOINT_FOR_DOC = "Qwen/Qwen2-7B-beta"
+_CONFIG_FOR_DOC = "Qwen2Config"
+
+QWEN2_PRETRAINED_MODEL_ARCHIVE_LIST = [
+    "Qwen/Qwen2-7B-beta",
+    # See all Qwen2 models at https://huggingface.co/models?filter=qwen2
+]
+
+from .mask_sdpa_utils import (
+    find_prefix_seq_length_by_pe,
+    update_causal_mask_with_pad_non_visible_2d,
+    update_causal_mask_for_one_gen_window_2d,
+    create_block_diff_mask_by_pe_4d,
+    find_pred_pos_from_input_ids
+)
+
+from .mask_magi_utils import build_magi_ranges
+
+# Copied from transformers.models.llama.modeling_llama._get_unpad_data
+def _get_unpad_data(attention_mask):
+    seqlens_in_batch = attention_mask.sum(dim=-1, dtype=torch.int32)
+    indices = torch.nonzero(attention_mask.flatten(), as_tuple=False).flatten()
+    max_seqlen_in_batch = seqlens_in_batch.max().item()
+    cu_seqlens = F.pad(torch.cumsum(seqlens_in_batch, dim=0, dtype=torch.torch.int32), (1, 0))
+    return (
+        indices,
+        cu_seqlens,
+        max_seqlen_in_batch,
+    )
+
+
+# Copied from transformers.models.llama.modeling_llama.LlamaRMSNorm with Llama->Qwen2
+class Qwen2RMSNorm(nn.Module):
+    def __init__(self, hidden_size, eps=1e-6):
+        """
+        Qwen2RMSNorm is equivalent to T5LayerNorm
+        """
+        super().__init__()
+        self.weight = nn.Parameter(torch.ones(hidden_size))
+        self.variance_epsilon = eps
+
+    def forward(self, hidden_states):
+        input_dtype = hidden_states.dtype
+        hidden_states = hidden_states.to(torch.float32)
+        variance = hidden_states.pow(2).mean(-1, keepdim=True)
+        hidden_states = hidden_states * torch.rsqrt(variance + self.variance_epsilon)
+        return self.weight * hidden_states.to(input_dtype)
+
+
+# Copied from transformers.models.llama.modeling_llama.LlamaRotaryEmbedding with Llama->Qwen2
+class Qwen2RotaryEmbedding(nn.Module):
+    def __init__(self, dim, max_position_embeddings=2048, base=10000, device=None):
+        super().__init__()
+
+        self.dim = dim
+        self.max_position_embeddings = max_position_embeddings
+        self.base = base
+        inv_freq = 1.0 / (self.base ** (torch.arange(0, self.dim, 2).float().to(device) / self.dim))
+        self.register_buffer("inv_freq", inv_freq, persistent=False)
+
+        # Build here to make `torch.jit.trace` work.
+        self._set_cos_sin_cache(
+            seq_len=max_position_embeddings, device=self.inv_freq.device, dtype=torch.get_default_dtype()
+        )
+
+    def _set_cos_sin_cache(self, seq_len, device, dtype):
+        self.max_seq_len_cached = seq_len
+        t = torch.arange(self.max_seq_len_cached, device=device, dtype=self.inv_freq.dtype)
+
+        freqs = torch.outer(t, self.inv_freq)
+        # Different from paper, but it uses a different permutation in order to obtain the same calculation
+        emb = torch.cat((freqs, freqs), dim=-1)
+        self.register_buffer("cos_cached", emb.cos().to(dtype), persistent=False)
+        self.register_buffer("sin_cached", emb.sin().to(dtype), persistent=False)
+
+    def forward(self, x, seq_len=None):
+        # x: [bs, num_attention_heads, seq_len, head_size]
+        if seq_len > self.max_seq_len_cached:
+            self._set_cos_sin_cache(seq_len=seq_len, device=x.device, dtype=x.dtype)
+
+        return (
+            self.cos_cached[:seq_len].to(dtype=x.dtype),
+            self.sin_cached[:seq_len].to(dtype=x.dtype),
+        )
+
+
+# Copied from transformers.models.llama.modeling_llama.rotate_half
+def rotate_half(x):
+    """Rotates half the hidden dims of the input."""
+    x1 = x[..., : x.shape[-1] // 2]
+    x2 = x[..., x.shape[-1] // 2 :]
+    return torch.cat((-x2, x1), dim=-1)
+
+
+# Copied from transformers.models.llama.modeling_llama.apply_rotary_pos_emb
+def apply_rotary_pos_emb(q, k, cos, sin, position_ids, unsqueeze_dim=1):
+    """Applies Rotary Position Embedding to the query and key tensors.
+
+    Args:
+        q (`torch.Tensor`): The query tensor.
+        k (`torch.Tensor`): The key tensor.
+        cos (`torch.Tensor`): The cosine part of the rotary embedding.
+        sin (`torch.Tensor`): The sine part of the rotary embedding.
+        position_ids (`torch.Tensor`):
+            The position indices of the tokens corresponding to the query and key tensors. For example, this can be
+            used to pass offsetted position ids when working with a KV-cache.
+        unsqueeze_dim (`int`, *optional*, defaults to 1):
+            The 'unsqueeze_dim' argument specifies the dimension along which to unsqueeze cos[position_ids] and
+            sin[position_ids] so that they can be properly broadcasted to the dimensions of q and k. For example, note
+            that cos[position_ids] and sin[position_ids] have the shape [batch_size, seq_len, head_dim]. Then, if q and
+            k have the shape [batch_size, heads, seq_len, head_dim], then setting unsqueeze_dim=1 makes
+            cos[position_ids] and sin[position_ids] broadcastable to the shapes of q and k. Similarly, if q and k have
+            the shape [batch_size, seq_len, heads, head_dim], then set unsqueeze_dim=2.
+    Returns:
+        `tuple(torch.Tensor)` comprising of the query and key tensors rotated using the Rotary Position Embedding.
+    """
+    cos = cos[position_ids].unsqueeze(unsqueeze_dim)
+    sin = sin[position_ids].unsqueeze(unsqueeze_dim)
+    q_embed = (q * cos) + (rotate_half(q) * sin)
+    k_embed = (k * cos) + (rotate_half(k) * sin)
+    return q_embed, k_embed
+
+
+# Copied from transformers.models.mistral.modeling_mistral.MistralMLP with Mistral->Qwen2
+class Qwen2MLP(nn.Module):
+    def __init__(self, config):
+        super().__init__()
+        self.config = config
+        self.hidden_size = config.hidden_size
+        self.intermediate_size = config.intermediate_size
+        self.gate_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
+        self.up_proj = nn.Linear(self.hidden_size, self.intermediate_size, bias=False)
+        self.down_proj = nn.Linear(self.intermediate_size, self.hidden_size, bias=False)
+        self.act_fn = ACT2FN[config.hidden_act]
+
+    def forward(self, x):
+        return self.down_proj(self.act_fn(self.gate_proj(x)) * self.up_proj(x))
+
+
+# Copied from transformers.models.llama.modeling_llama.repeat_kv
+def repeat_kv(hidden_states: torch.Tensor, n_rep: int) -> torch.Tensor:
+    """
+    This is the equivalent of torch.repeat_interleave(x, dim=1, repeats=n_rep). The hidden states go from (batch,
+    num_key_value_heads, seqlen, head_dim) to (batch, num_attention_heads, seqlen, head_dim)
+    """
+    batch, num_key_value_heads, slen, head_dim = hidden_states.shape
+    if n_rep == 1:
+        return hidden_states
+    hidden_states = hidden_states[:, :, None, :, :].expand(batch, num_key_value_heads, n_rep, slen, head_dim)
+    return hidden_states.reshape(batch, num_key_value_heads * n_rep, slen, head_dim)
+
+
+class Qwen2Attention(nn.Module):
+    """
+    Multi-headed attention from 'Attention Is All You Need' paper. Modified to use sliding window attention: Longformer
+    and "Generating Long Sequences with Sparse Transformers".
+    """
+
+    def __init__(self, config: Qwen2Config, layer_idx: Optional[int] = None):
+        super().__init__()
+        self.config = config
+        self.layer_idx = layer_idx
+        if layer_idx is None:
+            logger.warning_once(
+                f"Instantiating {self.__class__.__name__} without passing `layer_idx` is not recommended and will "
+                "to errors during the forward call, if caching is used. Please make sure to provide a `layer_idx` "
+                "when creating this class."
+            )
+
+        self.hidden_size = config.hidden_size
+        self.num_heads = config.num_attention_heads
+        self.head_dim = self.hidden_size // self.num_heads
+        self.num_key_value_heads = config.num_key_value_heads
+        self.num_key_value_groups = self.num_heads // self.num_key_value_heads
+        self.max_position_embeddings = config.max_position_embeddings
+        self.rope_theta = config.rope_theta
+        self.is_causal = True
+        self.attention_dropout = config.attention_dropout
+
+        if (self.head_dim * self.num_heads) != self.hidden_size:
+            raise ValueError(
+                f"hidden_size must be divisible by num_heads (got `hidden_size`: {self.hidden_size}"
+                f" and `num_heads`: {self.num_heads})."
+            )
+        self.q_proj = nn.Linear(self.hidden_size, self.num_heads * self.head_dim, bias=True)
+        self.k_proj = nn.Linear(self.hidden_size, self.num_key_value_heads * self.head_dim, bias=True)
+        self.v_proj = nn.Linear(self.hidden_size, self.num_key_value_heads * self.head_dim, bias=True)
+        self.o_proj = nn.Linear(self.num_heads * self.head_dim, self.hidden_size, bias=False)
+
+        self.rotary_emb = Qwen2RotaryEmbedding(
+            self.head_dim,
+            max_position_embeddings=self.max_position_embeddings,
+            base=self.rope_theta,
+        )
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: bool = False,
+        use_cache: bool = False,
+        **kwargs,
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
+        if "padding_mask" in kwargs:
+            warnings.warn(
+                "Passing `padding_mask` is deprecated and will be removed in v4.37. Please make sure use `attention_mask` instead.`"
+            )
+        bsz, q_len, _ = hidden_states.size()
+
+        query_states = self.q_proj(hidden_states)
+        key_states = self.k_proj(hidden_states)
+        value_states = self.v_proj(hidden_states)
+
+        query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+        key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+        value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+        kv_seq_len = key_states.shape[-2]
+        if past_key_value is not None:
+            if self.layer_idx is None:
+                raise ValueError(
+                    f"The cache structure has changed since version v4.36. If you are using {self.__class__.__name__} "
+                    "for auto-regressive decoding with k/v caching, please make sure to initialize the attention class "
+                    "with a layer index."
+                )
+            kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+        cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
+        query_states, key_states = apply_rotary_pos_emb(query_states, key_states, cos, sin, position_ids)
+
+        if past_key_value is not None:
+            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
+            key_states, value_states = past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs)
+
+        # repeat k/v heads if n_kv_heads < n_heads
+        key_states = repeat_kv(key_states, self.num_key_value_groups)
+        value_states = repeat_kv(value_states, self.num_key_value_groups)
+
+        attn_weights = torch.matmul(query_states, key_states.transpose(2, 3)) / math.sqrt(self.head_dim)
+
+        if attn_weights.size() != (bsz, self.num_heads, q_len, kv_seq_len):
+            raise ValueError(
+                f"Attention weights should be of size {(bsz, self.num_heads, q_len, kv_seq_len)}, but is"
+                f" {attn_weights.size()}"
+            )
+
+        if attention_mask is not None:
+            if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+                raise ValueError(
+                    f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, but is {attention_mask.size()}"
+                )
+
+            attn_weights = attn_weights + attention_mask
+
+        # upcast attention to fp32
+        attn_weights = nn.functional.softmax(attn_weights, dim=-1, dtype=torch.float32).to(query_states.dtype)
+        attn_weights = nn.functional.dropout(attn_weights, p=self.attention_dropout, training=self.training)
+        attn_output = torch.matmul(attn_weights, value_states)
+
+        if attn_output.size() != (bsz, self.num_heads, q_len, self.head_dim):
+            raise ValueError(
+                f"`attn_output` should be of size {(bsz, self.num_heads, q_len, self.head_dim)}, but is"
+                f" {attn_output.size()}"
+            )
+
+        attn_output = attn_output.transpose(1, 2).contiguous()
+        attn_output = attn_output.reshape(bsz, q_len, self.hidden_size)
+
+        attn_output = self.o_proj(attn_output)
+
+        if not output_attentions:
+            attn_weights = None
+
+        return attn_output, attn_weights, past_key_value
+
+
+class Qwen2FlashAttention2(Qwen2Attention):
+    """
+    Qwen2 flash attention module, following Qwen2 attention module. This module inherits from `Qwen2Attention`
+    as the weights of the module stays untouched. The only required change would be on the forward pass
+    where it needs to correctly call the public API of flash attention and deal with padding tokens
+    in case the input contains any of them. Additionally, for sliding window attention, we apply SWA only to the bottom
+    config.max_window_layers layers.
+    """
+
+    # Copied from transformers.models.llama.modeling_llama.LlamaFlashAttention2.__init__
+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+
+        # TODO: Should be removed once Flash Attention for RoCm is bumped to 2.1.
+        # flash_attn<2.1 generates top-left aligned causal mask, while what is needed here is bottom-right alignement, that was made default for flash_attn>=2.1. This attribute is used to handle this difference. Reference: https://github.com/Dao-AILab/flash-attention/releases/tag/v2.1.0.
+        # Beware that with flash_attn<2.1, using q_seqlen != k_seqlen (except for the case q_seqlen == 1) produces a wrong mask (top-left).
+        self._flash_attn_uses_top_left_mask = not is_flash_attn_greater_or_equal_2_10()
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: bool = False,
+        use_cache: bool = False,
+        **kwargs,
+    ):
+        if "padding_mask" in kwargs:
+            warnings.warn(
+                "Passing `padding_mask` is deprecated and will be removed in v4.37. Please make sure use `attention_mask` instead.`"
+            )
+
+            # overwrite attention_mask with padding_mask
+            attention_mask = kwargs.pop("padding_mask")
+        bsz, q_len, _ = hidden_states.size()
+
+        query_states = self.q_proj(hidden_states)
+        key_states = self.k_proj(hidden_states)
+        value_states = self.v_proj(hidden_states)
+
+        query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+        key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+        value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+        kv_seq_len = key_states.shape[-2]
+        if past_key_value is not None:
+            if self.layer_idx is None:
+                raise ValueError(
+                    f"The cache structure has changed since version v4.36. If you are using {self.__class__.__name__} "
+                    "for auto-regressive decoding with k/v caching, please make sure to initialize the attention class "
+                    "with a layer index."
+                )
+            kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+
+        # Because the input can be padded, the absolute sequence length depends on the max position id.
+        rotary_seq_len = max(kv_seq_len, position_ids[:, -1].max().item()) + 1
+        cos, sin = self.rotary_emb(value_states, seq_len=rotary_seq_len)
+
+        query_states, key_states = apply_rotary_pos_emb(query_states, key_states, cos, sin, position_ids)
+
+        use_sliding_windows = (
+            _flash_supports_window_size
+            and getattr(self.config, "sliding_window", None) is not None
+            and kv_seq_len > self.config.sliding_window
+            and self.config.use_sliding_window
+        )
+
+        if not _flash_supports_window_size:
+            logger.warning_once(
+                "The current flash attention version does not support sliding window attention, for a more memory efficient implementation"
+                " make sure to upgrade flash-attn library."
+            )
+
+        if past_key_value is not None:
+            # Activate slicing cache only if the config has a value `sliding_windows` attribute
+            cache_has_contents = past_key_value.get_seq_length(self.layer_idx) > 0
+            if (
+                getattr(self.config, "sliding_window", None) is not None
+                and kv_seq_len > self.config.sliding_window
+                and cache_has_contents
+            ):
+                slicing_tokens = 1 - self.config.sliding_window
+
+                past_key = past_key_value[self.layer_idx][0]
+                past_value = past_key_value[self.layer_idx][1]
+
+                past_key = past_key[:, :, slicing_tokens:, :].contiguous()
+                past_value = past_value[:, :, slicing_tokens:, :].contiguous()
+
+                if past_key.shape[-2] != self.config.sliding_window - 1:
+                    raise ValueError(
+                        f"past key must have a shape of (`batch_size, num_heads, self.config.sliding_window-1, head_dim`), got"
+                        f" {past_key.shape}"
+                    )
+
+                if attention_mask is not None:
+                    attention_mask = attention_mask[:, slicing_tokens:]
+                    attention_mask = torch.cat([attention_mask, torch.ones_like(attention_mask[:, -1:])], dim=-1)
+
+            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
+            key_states, value_states = past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs)
+
+        # repeat k/v heads if n_kv_heads < n_heads
+        key_states = repeat_kv(key_states, self.num_key_value_groups)
+        value_states = repeat_kv(value_states, self.num_key_value_groups)
+        dropout_rate = 0.0 if not self.training else self.attention_dropout
+
+        # In PEFT, usually we cast the layer norms in float32 for training stability reasons
+        # therefore the input hidden states gets silently casted in float32. Hence, we need
+        # cast them back in float16 just to be sure everything works as expected.
+        input_dtype = query_states.dtype
+        if input_dtype == torch.float32:
+            if torch.is_autocast_enabled():
+                target_dtype = torch.get_autocast_gpu_dtype()
+            # Handle the case where the model is quantized
+            elif hasattr(self.config, "_pre_quantization_dtype"):
+                target_dtype = self.config._pre_quantization_dtype
+            else:
+                target_dtype = self.q_proj.weight.dtype
+
+            logger.warning_once(
+                f"The input hidden states seems to be silently casted in float32, this might be related to"
+                f" the fact you have upcasted embedding or layer norm layers in float32. We will cast back the input in"
+                f" {target_dtype}."
+            )
+
+            query_states = query_states.to(target_dtype)
+            key_states = key_states.to(target_dtype)
+            value_states = value_states.to(target_dtype)
+
+        # Reashape to the expected shape for Flash Attention
+        query_states = query_states.transpose(1, 2)
+        key_states = key_states.transpose(1, 2)
+        value_states = value_states.transpose(1, 2)
+
+        attn_output = self._flash_attention_forward(
+            query_states,
+            key_states,
+            value_states,
+            attention_mask,
+            q_len,
+            dropout=dropout_rate,
+            use_sliding_windows=use_sliding_windows,
+        )
+
+        attn_output = attn_output.reshape(bsz, q_len, self.hidden_size).contiguous()
+        attn_output = self.o_proj(attn_output)
+
+        if not output_attentions:
+            attn_weights = None
+
+        return attn_output, attn_weights, past_key_value
+
+    def _flash_attention_forward(
+        self,
+        query_states,
+        key_states,
+        value_states,
+        attention_mask,
+        query_length,
+        dropout=0.0,
+        softmax_scale=None,
+        use_sliding_windows=False,
+    ):
+        """
+        Calls the forward method of Flash Attention - if the input hidden states contain at least one padding token
+        first unpad the input, then computes the attention scores and pad the final attention scores.
+
+        Args:
+            query_states (`torch.Tensor`):
+                Input query states to be passed to Flash Attention API
+            key_states (`torch.Tensor`):
+                Input key states to be passed to Flash Attention API
+            value_states (`torch.Tensor`):
+                Input value states to be passed to Flash Attention API
+            attention_mask (`torch.Tensor`):
+                The padding mask - corresponds to a tensor of size `(batch_size, seq_len)` where 0 stands for the
+                position of padding tokens and 1 for the position of non-padding tokens.
+            dropout (`int`, *optional*):
+                Attention dropout
+            softmax_scale (`float`, *optional*):
+                The scaling of QK^T before applying softmax. Default to 1 / sqrt(head_dim)
+            use_sliding_windows (`bool`, *optional*):
+                Whether to activate sliding window attention.
+        """
+        if not self._flash_attn_uses_top_left_mask:
+            causal = self.is_causal
+        else:
+            # TODO: Remove the `query_length != 1` check once Flash Attention for RoCm is bumped to 2.1. For details, please see the comment in LlamaFlashAttention2 __init__.
+            causal = self.is_causal and query_length != 1
+
+        # Decide whether to use SWA or not by layer index.
+        if use_sliding_windows and self.layer_idx >= self.config.max_window_layers:
+            use_sliding_windows = False
+
+        # Contains at least one padding token in the sequence
+        if attention_mask is not None:
+            batch_size = query_states.shape[0]
+            query_states, key_states, value_states, indices_q, cu_seq_lens, max_seq_lens = self._upad_input(
+                query_states, key_states, value_states, attention_mask, query_length
+            )
+
+            cu_seqlens_q, cu_seqlens_k = cu_seq_lens
+            max_seqlen_in_batch_q, max_seqlen_in_batch_k = max_seq_lens
+
+            if not use_sliding_windows:
+                attn_output_unpad = flash_attn_varlen_func(
+                    query_states,
+                    key_states,
+                    value_states,
+                    cu_seqlens_q=cu_seqlens_q,
+                    cu_seqlens_k=cu_seqlens_k,
+                    max_seqlen_q=max_seqlen_in_batch_q,
+                    max_seqlen_k=max_seqlen_in_batch_k,
+                    dropout_p=dropout,
+                    softmax_scale=softmax_scale,
+                    causal=causal,
+                )
+            else:
+                attn_output_unpad = flash_attn_varlen_func(
+                    query_states,
+                    key_states,
+                    value_states,
+                    cu_seqlens_q=cu_seqlens_q,
+                    cu_seqlens_k=cu_seqlens_k,
+                    max_seqlen_q=max_seqlen_in_batch_q,
+                    max_seqlen_k=max_seqlen_in_batch_k,
+                    dropout_p=dropout,
+                    softmax_scale=softmax_scale,
+                    causal=causal,
+                    window_size=(self.config.sliding_window, self.config.sliding_window),
+                )
+
+            attn_output = pad_input(attn_output_unpad, indices_q, batch_size, query_length)
+        else:
+            if not use_sliding_windows:
+                attn_output = flash_attn_func(
+                    query_states,
+                    key_states,
+                    value_states,
+                    dropout,
+                    softmax_scale=softmax_scale,
+                    causal=causal,
+                )
+            else:
+                attn_output = flash_attn_func(
+                    query_states,
+                    key_states,
+                    value_states,
+                    dropout,
+                    softmax_scale=softmax_scale,
+                    causal=causal,
+                    window_size=(self.config.sliding_window, self.config.sliding_window),
+                )
+
+        return attn_output
+
+    # Copied from transformers.models.mistral.modeling_mistral.MistralFlashAttention2._upad_input
+    def _upad_input(self, query_layer, key_layer, value_layer, attention_mask, query_length):
+        batch_size, kv_seq_len, num_heads, head_dim = key_layer.shape
+
+        # On the first iteration we need to properly re-create the padding mask
+        # by slicing it on the proper place
+        if kv_seq_len != attention_mask.shape[-1]:
+            attention_mask_num_tokens = attention_mask.shape[-1]
+            attention_mask = attention_mask[:, attention_mask_num_tokens - kv_seq_len :]
+
+        indices_k, cu_seqlens_k, max_seqlen_in_batch_k = _get_unpad_data(attention_mask)
+
+        key_layer = index_first_axis(key_layer.reshape(batch_size * kv_seq_len, num_heads, head_dim), indices_k)
+        value_layer = index_first_axis(value_layer.reshape(batch_size * kv_seq_len, num_heads, head_dim), indices_k)
+
+        if query_length == kv_seq_len:
+            query_layer = index_first_axis(
+                query_layer.reshape(batch_size * kv_seq_len, num_heads, head_dim), indices_k
+            )
+            cu_seqlens_q = cu_seqlens_k
+            max_seqlen_in_batch_q = max_seqlen_in_batch_k
+            indices_q = indices_k
+        elif query_length == 1:
+            max_seqlen_in_batch_q = 1
+            cu_seqlens_q = torch.arange(
+                batch_size + 1, dtype=torch.int32, device=query_layer.device
+            )  # There is a memcpy here, that is very bad.
+            indices_q = cu_seqlens_q[:-1]
+            query_layer = query_layer.squeeze(1)
+        else:
+            # The -q_len: slice assumes left padding.
+            attention_mask = attention_mask[:, -query_length:]
+            query_layer, indices_q, cu_seqlens_q, max_seqlen_in_batch_q = unpad_input(query_layer, attention_mask)
+
+        return (
+            query_layer,
+            key_layer,
+            value_layer,
+            indices_q,
+            (cu_seqlens_q, cu_seqlens_k),
+            (max_seqlen_in_batch_q, max_seqlen_in_batch_k),
+        )
+
+
+# Copied from transformers.models.llama.modeling_llama.LlamaSdpaAttention with Llama->Qwen2
+class Qwen2SdpaAttention(Qwen2Attention):
+    """
+    Qwen2 attention module using torch.nn.functional.scaled_dot_product_attention. This module inherits from
+    `Qwen2Attention` as the weights of the module stays untouched. The only changes are on the forward pass to adapt to
+    SDPA API.
+    """
+
+    # Adapted from Qwen2Attention.forward
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: bool = False,
+        use_cache: bool = False,
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
+        if output_attentions:
+            # TODO: Improve this warning with e.g. `model.config.attn_implementation = "manual"` once this is implemented.
+            logger.warning_once(
+                "Qwen2Model is using Qwen2SdpaAttention, but `torch.nn.functional.scaled_dot_product_attention` does not support `output_attentions=True`. Falling back to the manual attention implementation, "
+                'but specifying the manual implementation will be required from Transformers version v5.0.0 onwards. This warning can be removed using the argument `attn_implementation="eager"` when loading the model.'
+            )
+            return super().forward(
+                hidden_states=hidden_states,
+                attention_mask=attention_mask,
+                position_ids=position_ids,
+                past_key_value=past_key_value,
+                output_attentions=output_attentions,
+                use_cache=use_cache,
+            )
+
+        bsz, q_len, _ = hidden_states.size()
+
+        query_states = self.q_proj(hidden_states)
+        key_states = self.k_proj(hidden_states)
+        value_states = self.v_proj(hidden_states)
+
+        query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+        key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+        value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+        kv_seq_len = key_states.shape[-2]
+        if past_key_value is not None:
+            kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+        cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
+
+        query_states, key_states = apply_rotary_pos_emb(query_states, key_states, cos, sin, position_ids)
+
+        if past_key_value is not None:
+            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
+            key_states, value_states = past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs)
+
+        key_states = repeat_kv(key_states, self.num_key_value_groups)
+        value_states = repeat_kv(value_states, self.num_key_value_groups)
+
+        if attention_mask is not None:
+            if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+                raise ValueError(
+                    f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, but is {attention_mask.size()}"
+                )
+
+        # SDPA with memory-efficient backend is currently (torch==2.1.2) bugged with non-contiguous inputs with custom attn_mask,
+        # Reference: https://github.com/pytorch/pytorch/issues/112577.
+        if query_states.device.type == "cuda" and attention_mask is not None:
+            query_states = query_states.contiguous()
+            key_states = key_states.contiguous()
+            value_states = value_states.contiguous()
+
+        attn_output = torch.nn.functional.scaled_dot_product_attention(
+            query_states,
+            key_states,
+            value_states,
+            attn_mask=attention_mask,
+            dropout_p=self.attention_dropout if self.training else 0.0,
+            is_causal=False,
+        )
+
+        attn_output = attn_output.transpose(1, 2).contiguous()
+        attn_output = attn_output.reshape(bsz, q_len, self.hidden_size)
+
+        attn_output = self.o_proj(attn_output)
+
+        return attn_output, None, past_key_value
+
+
+class Qwen2SdpaAttentionGqa(Qwen2Attention):
+    """
+    Qwen2 attention module using torch.nn.functional.scaled_dot_product_attention. This module inherits from
+    `Qwen2Attention` as the weights of the module stays untouched. The only changes are on the forward pass to adapt to
+    SDPA API.
+    """
+
+    # Adapted from Qwen2Attention.forward
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: bool = False,
+        use_cache: bool = False,
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
+        if output_attentions:
+            # TODO: Improve this warning with e.g. `model.config.attn_implementation = "manual"` once this is implemented.
+            logger.warning_once(
+                "Qwen2Model is using Qwen2SdpaAttention, but `torch.nn.functional.scaled_dot_product_attention` does not support `output_attentions=True`. Falling back to the manual attention implementation, "
+                'but specifying the manual implementation will be required from Transformers version v5.0.0 onwards. This warning can be removed using the argument `attn_implementation="eager"` when loading the model.'
+            )
+            return super().forward(
+                hidden_states=hidden_states,
+                attention_mask=attention_mask,
+                position_ids=position_ids,
+                past_key_value=past_key_value,
+                output_attentions=output_attentions,
+                use_cache=use_cache,
+            )
+
+        bsz, q_len, _ = hidden_states.size()
+
+        query_states = self.q_proj(hidden_states)
+        key_states = self.k_proj(hidden_states)
+        value_states = self.v_proj(hidden_states)
+
+        query_states = query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2)
+        key_states = key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+        value_states = value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2)
+
+        kv_seq_len = key_states.shape[-2]
+        if past_key_value is not None:
+            kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+        cos, sin = self.rotary_emb(value_states, seq_len=kv_seq_len)
+
+        query_states, key_states = apply_rotary_pos_emb(query_states, key_states, cos, sin, position_ids)
+
+        if past_key_value is not None:
+            cache_kwargs = {"sin": sin, "cos": cos}  # Specific to RoPE models
+            key_states, value_states = past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs)
+
+        # key_states = repeat_kv(key_states, self.num_key_value_groups)
+        # value_states = repeat_kv(value_states, self.num_key_value_groups)
+
+        if attention_mask is not None:
+            if attention_mask.size() != (bsz, 1, q_len, kv_seq_len):
+                raise ValueError(
+                    f"Attention mask should be of size {(bsz, 1, q_len, kv_seq_len)}, but is {attention_mask.size()}"
+                )
+
+        # SDPA with memory-efficient backend is currently (torch==2.1.2) bugged with non-contiguous inputs with custom attn_mask,
+        # Reference: https://github.com/pytorch/pytorch/issues/112577.
+        if query_states.device.type == "cuda" and attention_mask is not None:
+            query_states = query_states.contiguous()
+            key_states = key_states.contiguous()
+            value_states = value_states.contiguous()
+
+        with torch.backends.cuda.sdp_kernel(enable_flash=True,   
+                                    enable_math=True, 
+                                    enable_mem_efficient=False):
+
+            attn_output = torch.nn.functional.scaled_dot_product_attention(
+                query_states,
+                key_states,
+                value_states,
+                attn_mask=attention_mask,
+                enable_gqa=True,
+                dropout_p=self.attention_dropout if self.training else 0.0,
+                is_causal=False,
+            )
+
+        attn_output = attn_output.transpose(1, 2).contiguous()
+        attn_output = attn_output.reshape(bsz, q_len, self.hidden_size)
+
+        attn_output = self.o_proj(attn_output)
+
+        return attn_output, None, past_key_value
+
+
+class Qwen2MagiAttention(Qwen2Attention):
+    """
+    Qwen2 attention using MagiAttention for efficient training with MTP packing support.
+    
+    MagiAttention uses range-based sparse attention patterns:
+    - q_ranges/k_ranges define which query/key ranges attend to each other
+    - attn_type_map specifies causal(1) or full(0) attention for each range pair
+    """
+    
+    def __init__(self, *args, **kwargs):
+        super().__init__(*args, **kwargs)
+        if not _MAGI_AVAILABLE:
+            raise ImportError(
+                "magi_attention is not installed. Install with: pip install magi-attention"
+            )
+        self.softmax_scale = self.head_dim ** -0.5
+    
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[dict] = None,  # magi_plan dict
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Cache] = None,
+        output_attentions: bool = False,
+        use_cache: bool = False,
+    ) -> Tuple[torch.Tensor, Optional[torch.Tensor], Optional[Tuple[torch.Tensor]]]:
+        if output_attentions:
+            raise NotImplementedError('MagiAttention does not support output_attentions=True')
+
+        bsz, q_len, _ = hidden_states.size()
+        assert bsz == 1, "MagiAttention only supports batch_size=1 (use packing instead)"
+
+        query_states = self.q_proj(hidden_states)
+        key_states = self.k_proj(hidden_states)
+        value_states = self.v_proj(hidden_states)
+
+        # Magi expects [T, H, D] format (no batch dimension)
+        query_states = query_states.view(q_len, self.num_heads, self.head_dim)
+        key_states = key_states.view(q_len, self.num_key_value_heads, self.head_dim)
+        value_states = value_states.view(q_len, self.num_key_value_heads, self.head_dim)
+
+        kv_seq_len = q_len
+        if past_key_value is not None:
+            if self.layer_idx is None:
+                raise ValueError(
+                    f"The cache structure has changed since version v4.36. If you are using {self.__class__.__name__} "
+                    "for auto-regressive decoding with k/v caching, please make sure to initialize the attention class "
+                    "with a layer index."
+                )
+            kv_seq_len += past_key_value.get_seq_length(self.layer_idx)
+        
+        cos, sin = self.rotary_emb(value_states.unsqueeze(0).transpose(1, 2), seq_len=kv_seq_len)
+        
+        # Apply RoPE: need [B, H, L, D] format for apply_rotary_pos_emb
+        q_for_rope = query_states.unsqueeze(0).transpose(1, 2)  # [1, H, L, D]
+        k_for_rope = key_states.unsqueeze(0).transpose(1, 2)    # [1, Hkv, L, D]
+        q_for_rope, k_for_rope = apply_rotary_pos_emb(q_for_rope, k_for_rope, cos, sin, position_ids)
+        
+        # Back to [T, H, D]
+        query_states = q_for_rope.squeeze(0).transpose(0, 1).contiguous()  # [L, H, D]
+        key_states = k_for_rope.squeeze(0).transpose(0, 1).contiguous()    # [L, Hkv, D]
+
+        if past_key_value is not None:
+            cache_kwargs = {"sin": sin, "cos": cos}
+            # Note: Magi doesn't support KV cache in training, this is for potential future use
+            key_states_4d = key_states.unsqueeze(0).transpose(1, 2)
+            value_states_4d = value_states.unsqueeze(0).transpose(1, 2)
+            key_states_4d, value_states_4d = past_key_value.update(
+                key_states_4d, value_states_4d, self.layer_idx, cache_kwargs
+            )
+            key_states = key_states_4d.squeeze(0).transpose(0, 1).contiguous()
+            value_states = value_states_4d.squeeze(0).transpose(0, 1).contiguous()
+
+        # Run Magi Attention
+        # attention_mask is a magi_plan dict with q_ranges, k_ranges, attn_type_map, etc.
+
+        attn_output, _ = flex_flash_attn_func(
+            query_states.contiguous(),
+            key_states.contiguous(),
+            value_states.contiguous(),
+            q_ranges=attention_mask["q_ranges"],
+            k_ranges=attention_mask["k_ranges"],
+            attn_type_map=attention_mask["attn_type_map"],
+            softmax_scale=self.softmax_scale,
+            softcap=0.0,
+            deterministic=False,
+        )  # [T, H, D]
+        
+        # Reshape to [B, L, H*D]
+        attn_output = attn_output.view(1, q_len, self.hidden_size)
+        attn_output = self.o_proj(attn_output)
+
+        return attn_output, None, past_key_value
+
+
+QWEN2_ATTENTION_CLASSES = {
+    "eager": Qwen2Attention,
+    "flash_attention_2": Qwen2FlashAttention2,
+    "sdpa": Qwen2SdpaAttention,
+    "magi": Qwen2MagiAttention,
+}
+
+
+class Qwen2DecoderLayer(nn.Module):
+    def __init__(self, config: Qwen2Config, layer_idx: int):
+        super().__init__()
+        self.hidden_size = config.hidden_size
+
+        if config._attn_implementation == 'magi' and not _MAGI_AVAILABLE:
+            if is_flash_attn_2_available():
+                logger.warning_once(
+                    'magi_attention not available, falling back to flash_attention_2'
+                )
+                config._attn_implementation = 'flash_attention_2'
+            else:
+                logger.warning_once(
+                    'magi_attention not available, falling back to sdpa'
+                )
+                config._attn_implementation = 'sdpa'
+        if config._attn_implementation == 'flash_attention_2' and not is_flash_attn_2_available():
+            logger.warning_once(
+                'flash_attn is not available, falling back to sdpa'
+            )
+            config._attn_implementation = 'sdpa'
+
+        self.self_attn = QWEN2_ATTENTION_CLASSES[config._attn_implementation](config, layer_idx)
+
+        self.mlp = Qwen2MLP(config)
+        self.input_layernorm = Qwen2RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+        self.post_attention_layernorm = Qwen2RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_value: Optional[Tuple[torch.Tensor]] = None,
+        output_attentions: Optional[bool] = False,
+        use_cache: Optional[bool] = False,
+        **kwargs,
+    ) -> Tuple[torch.FloatTensor, Optional[Tuple[torch.FloatTensor, torch.FloatTensor]]]:
+        if "padding_mask" in kwargs:
+            warnings.warn(
+                "Passing `padding_mask` is deprecated and will be removed in v4.37. "
+                "Please make sure use `attention_mask` instead.`"
+            )
+        """
+        Args:
+            hidden_states (`torch.FloatTensor`): input to the layer of shape `(batch, seq_len, embed_dim)`
+            attention_mask (`torch.FloatTensor`, *optional*): attention mask of size
+                `(batch, sequence_length)` where padding elements are indicated by 0.
+            output_attentions (`bool`, *optional*):
+                Whether or not to return the attentions tensors of all attention layers. See `attentions` under
+                returned tensors for more detail.
+            use_cache (`bool`, *optional*):
+                If set to `True`, `past_key_values` key value states are returned and can be used to speed up decoding
+                (see `past_key_values`).
+            past_key_value (`Tuple(torch.FloatTensor)`, *optional*): cached past key and value projection states
+        """
+
+        residual = hidden_states
+
+        hidden_states = self.input_layernorm(hidden_states)
+
+        # Self Attention
+        hidden_states, self_attn_weights, present_key_value = self.self_attn(
+            hidden_states=hidden_states,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_value=past_key_value,
+            output_attentions=output_attentions,
+            use_cache=use_cache,
+        )
+        hidden_states = residual + hidden_states
+
+        # Fully Connected
+        residual = hidden_states
+        hidden_states = self.post_attention_layernorm(hidden_states)
+        hidden_states = self.mlp(hidden_states)
+        hidden_states = residual + hidden_states
+
+        outputs = (hidden_states,)
+
+        if output_attentions:
+            outputs += (self_attn_weights,)
+
+        if use_cache:
+            outputs += (present_key_value,)
+
+        return outputs
+
+
+QWEN2_START_DOCSTRING = r"""
+    This model inherits from [`PreTrainedModel`]. Check the superclass documentation for the generic methods the
+    library implements for all its model (such as downloading or saving, resizing the input embeddings, pruning heads
+    etc.)
+
+    This model is also a PyTorch [torch.nn.Module](https://pytorch.org/docs/stable/nn.html#torch.nn.Module) subclass.
+    Use it as a regular PyTorch Module and refer to the PyTorch documentation for all matter related to general usage
+    and behavior.
+
+    Parameters:
+        config ([`Qwen2Config`]):
+            Model configuration class with all the parameters of the model. Initializing with a config file does not
+            load the weights associated with the model, only the configuration. Check out the
+            [`~PreTrainedModel.from_pretrained`] method to load the model weights.
+"""
+
+
+@add_start_docstrings(
+    "The bare Qwen2 Model outputting raw hidden-states without any specific head on top.",
+    QWEN2_START_DOCSTRING,
+)
+class Qwen2PreTrainedModel(PreTrainedModel):
+    config_class = Qwen2Config
+    base_model_prefix = "model"
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["Qwen2DecoderLayer"]
+    _skip_keys_device_placement = "past_key_values"
+    _supports_flash_attn_2 = True
+    _supports_sdpa = True
+    _supports_cache_class = True
+
+    @classmethod
+    def _autoset_attn_implementation(cls, config, *args, **kwargs):
+        if getattr(config, '_attn_implementation', None) == 'magi':
+            return config
+        return super()._autoset_attn_implementation(config, *args, **kwargs)
+
+    def _check_and_adjust_attn_implementation(self, attn_implementation, is_init_check=False):
+        if attn_implementation == "magi":
+            return "magi"
+        return super()._check_and_adjust_attn_implementation(attn_implementation, is_init_check)
+
+    def _init_weights(self, module):
+        std = self.config.initializer_range
+        if isinstance(module, nn.Linear):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.bias is not None:
+                module.bias.data.zero_()
+        elif isinstance(module, nn.Embedding):
+            module.weight.data.normal_(mean=0.0, std=std)
+            if module.padding_idx is not None:
+                module.weight.data[module.padding_idx].zero_()
+
+
+QWEN2_INPUTS_DOCSTRING = r"""
+    Args:
+        input_ids (`torch.LongTensor` of shape `(batch_size, sequence_length)`):
+            Indices of input sequence tokens in the vocabulary. Padding will be ignored by default should you provide
+            it.
+
+            Indices can be obtained using [`AutoTokenizer`]. See [`PreTrainedTokenizer.encode`] and
+            [`PreTrainedTokenizer.__call__`] for details.
+
+            [What are input IDs?](../glossary#input-ids)
+        attention_mask (`torch.Tensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Mask to avoid performing attention on padding token indices. Mask values selected in `[0, 1]`:
+
+            - 1 for tokens that are **not masked**,
+            - 0 for tokens that are **masked**.
+
+            [What are attention masks?](../glossary#attention-mask)
+
+            Indices can be obtained using [`AutoTokenizer`]. See [`PreTrainedTokenizer.encode`] and
+            [`PreTrainedTokenizer.__call__`] for details.
+
+            If `past_key_values` is used, optionally only the last `decoder_input_ids` have to be input (see
+            `past_key_values`).
+
+            If you want to change padding behavior, you should read [`modeling_opt._prepare_decoder_attention_mask`]
+            and modify to your needs. See diagram 1 in [the paper](https://arxiv.org/abs/1910.13461) for more
+            information on the default strategy.
+
+            - 1 indicates the head is **not masked**,
+            - 0 indicates the head is **masked**.
+        position_ids (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Indices of positions of each input sequence tokens in the position embeddings. Selected in the range `[0,
+            config.n_positions - 1]`.
+
+            [What are position IDs?](../glossary#position-ids)
+        past_key_values (`Cache` or `tuple(tuple(torch.FloatTensor))`, *optional*):
+            Pre-computed hidden-states (key and values in the self-attention blocks and in the cross-attention
+            blocks) that can be used to speed up sequential decoding. This typically consists in the `past_key_values`
+            returned by the model at a previous stage of decoding, when `use_cache=True` or `config.use_cache=True`.
+
+            Two formats are allowed:
+            - a [`~cache_utils.Cache`] instance;
+            - Tuple of `tuple(torch.FloatTensor)` of length `config.n_layers`, with each tuple having 2 tensors of
+            shape `(batch_size, num_heads, sequence_length, embed_size_per_head)`). This is also known as the legacy
+            cache format.
+
+            The model will output the same cache format that is fed as input. If no `past_key_values` are passed, the
+            legacy cache format will be returned.
+
+            If `past_key_values` are used, the user can optionally input only the last `input_ids` (those that don't
+            have their past key value states given to this model) of shape `(batch_size, 1)` instead of all `input_ids`
+            of shape `(batch_size, sequence_length)`.
+        inputs_embeds (`torch.FloatTensor` of shape `(batch_size, sequence_length, hidden_size)`, *optional*):
+            Optionally, instead of passing `input_ids` you can choose to directly pass an embedded representation. This
+            is useful if you want more control over how to convert `input_ids` indices into associated vectors than the
+            model's internal embedding lookup matrix.
+        use_cache (`bool`, *optional*):
+            If set to `True`, `past_key_values` key value states are returned and can be used to speed up decoding (see
+            `past_key_values`).
+        output_attentions (`bool`, *optional*):
+            Whether or not to return the attentions tensors of all attention layers. See `attentions` under returned
+            tensors for more detail.
+        output_hidden_states (`bool`, *optional*):
+            Whether or not to return the hidden states of all layers. See `hidden_states` under returned tensors for
+            more detail.
+        return_dict (`bool`, *optional*):
+            Whether or not to return a [`~utils.ModelOutput`] instead of a plain tuple.
+"""
+
+
+@add_start_docstrings(
+    "The bare Qwen2 Model outputting raw hidden-states without any specific head on top.",
+    QWEN2_START_DOCSTRING,
+)
+class Qwen2Model(Qwen2PreTrainedModel):
+    """
+    Transformer decoder consisting of *config.num_hidden_layers* layers. Each layer is a [`Qwen2DecoderLayer`]
+
+    Args:
+        config: Qwen2Config
+    """
+
+    def __init__(self, config: Qwen2Config):
+        super().__init__(config)
+        self.padding_idx = config.pad_token_id
+        self.vocab_size = config.vocab_size
+
+        self.embed_tokens = nn.Embedding(config.vocab_size, config.hidden_size, self.padding_idx)
+        self.layers = nn.ModuleList(
+            [Qwen2DecoderLayer(config, layer_idx) for layer_idx in range(config.num_hidden_layers)]
+        )
+        self._attn_implementation = config._attn_implementation
+        self.norm = Qwen2RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
+
+        self.gradient_checkpointing = False
+        # Initialize weights and apply final processing
+        self.post_init()
+        
+        self.block_size = getattr(config, 'block_size', 6)
+        self.causal_attn = getattr(config, 'causal_attn', False)
+        self.text_mask_token_id = getattr(config, 'text_mask_token_id', 151676)
+
+
+    def get_input_embeddings(self):
+        return self.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.embed_tokens = value
+
+    def image_processing(self, input_ids, visual_features, image_token_index):
+        if visual_features is not None:
+            input_embeds = self.get_input_embeddings()(input_ids)
+            B, N, C = input_embeds.shape
+            input_embeds = input_embeds.reshape(B * N, C)
+
+            input_ids = input_ids.reshape(B * N)
+            selected = (input_ids == image_token_index)
+            assert selected.sum() != 0
+            input_embeds[selected] = visual_features.reshape(-1, C).to(input_embeds.device)
+            input_embeds = input_embeds.reshape(B, N, C)
+        else:
+            input_embeds = self.get_input_embeddings()(input_ids)
+        return input_embeds
+        
+    @add_start_docstrings_to_model_forward(QWEN2_INPUTS_DOCSTRING)
+    def forward(
+        self,
+        input_ids: torch.LongTensor = None,
+        visual_features: Optional[torch.FloatTensor] = None,
+        image_token_index: int = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[List[torch.FloatTensor]] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        use_cache: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        return_dict: Optional[bool] = None,
+    ) -> Union[Tuple, BaseModelOutputWithPast]:
+        output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
+        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
+
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        # retrieve input_ids and inputs_embeds
+        if input_ids is not None and inputs_embeds is not None:
+            raise ValueError("You cannot specify both decoder_input_ids and decoder_inputs_embeds at the same time")
+        elif input_ids is not None:
+            batch_size, seq_length = input_ids.shape
+        elif inputs_embeds is not None:
+            batch_size, seq_length, _ = inputs_embeds.shape
+        else:
+            raise ValueError("You have to specify either decoder_input_ids or decoder_inputs_embeds")
+
+        if self.gradient_checkpointing and self.training:
+            if use_cache:
+                logger.warning_once(
+                    "`use_cache=True` is incompatible with gradient checkpointing. Setting `use_cache=False`..."
+                )
+                use_cache = False
+
+        past_key_values_length = 0
+
+        if use_cache:
+            use_legacy_cache = not isinstance(past_key_values, Cache)
+            if use_legacy_cache:
+                if past_key_values is None:
+                    past_key_values = DynamicCache()
+                else:
+                    past_key_values = DynamicCache.from_legacy_cache(past_key_values)
+            past_key_values_length = past_key_values.get_seq_length()
+
+        if position_ids is None:
+            device = input_ids.device if input_ids is not None else inputs_embeds.device
+            position_ids = torch.arange(
+                past_key_values_length, seq_length + past_key_values_length, dtype=torch.long, device=device
+            )
+            position_ids = position_ids.unsqueeze(0).view(-1, seq_length)
+        else:
+            position_ids = position_ids.view(-1, seq_length).long()
+
+        if inputs_embeds is None:
+            inputs_embeds = self.image_processing(input_ids, visual_features, image_token_index)
+
+        if attention_mask is not None and self._attn_implementation == "magi" and use_cache:
+            is_padding_right = attention_mask[:, -1].sum().item() != batch_size
+            if is_padding_right:
+                raise ValueError(
+                    "You are attempting to perform batched generation with padding_side='right'"
+                    " this may lead to unexpected behaviour for Flash Attention version of Qwen2. Make sure to "
+                    " call `tokenizer.padding_side  = 'left'` before tokenizing the input. "
+                )
+
+        device = input_ids.device if input_ids is not None else inputs_embeds.device
+
+        x0_len = find_prefix_seq_length_by_pe(position_ids).to(device=device)
+
+        def _prepare_block_mask_for_inference(attention_mask):
+            attention_mask = _prepare_4d_causal_attention_mask(
+                attention_mask,
+                (batch_size, seq_length),
+                inputs_embeds,
+                past_key_values_length,
+                sliding_window=self.config.sliding_window,
+            )
+            # switch to ar mode
+            if seq_length == 1 or (input_ids is not None and input_ids[0][-1].item() != self.text_mask_token_id):
+                return attention_mask
+
+            
+            if attention_mask is None or len(attention_mask.shape) != 4:
+                return attention_mask
+
+            # For SDLM, the generation window should set to bidirectional attention
+            if use_cache:
+                update_mask_func = partial(
+                    update_causal_mask_for_one_gen_window_2d,
+                    block_size=self.block_size,
+                    use_cache=use_cache,
+                    causal_attn=self.causal_attn,
+                )
+            else:
+                update_mask_func = partial(
+                    update_causal_mask_with_pad_non_visible_2d,
+                    block_size=self.block_size,
+                    text_mask_token_id=self.text_mask_token_id,
+                    causal_attn=self.causal_attn,
+                )
+
+            new_attention_mask = []
+            for b in range(attention_mask.shape[0]):
+                new_attention_mask.append(
+                    update_mask_func(
+                        input_ids[b],
+                        attention_mask[b][0],
+                    ).unsqueeze(0)
+                )
+            return torch.stack(new_attention_mask, dim=0)
+
+        def _prepare_block_mask_for_training():
+            block_mask, _ = create_block_diff_mask_by_pe_4d(
+                block_size=self.block_size,
+                x0_len_list=x0_len,
+                position_ids=position_ids,
+                causal_attn=self.causal_attn,
+            )
+            return block_mask
+
+        if self._attn_implementation == "magi":
+            ar_decode = seq_length == 1 or (input_ids is not None and input_ids[0][-1].item() != self.text_mask_token_id)
+            attention_mask = build_magi_ranges(
+                kv_len=seq_length + past_key_values_length,
+                q_len=seq_length,
+                block_size=self.block_size,
+                ar_decode=ar_decode,
+                device=device
+            )
+
+        elif self._attn_implementation == "sdpa":
+            attention_mask = _prepare_block_mask_for_training() if self.training else _prepare_block_mask_for_inference(attention_mask)
+
+        else:
+            raise NotImplementedError(f'{self._attn_implementation=}')
+
+
+        hidden_states = inputs_embeds
+
+        # decoder layers
+        all_hidden_states = () if output_hidden_states else None
+        all_self_attns = () if output_attentions else None
+        next_decoder_cache = None
+
+        for decoder_layer in self.layers:
+            if output_hidden_states:
+                all_hidden_states += (hidden_states,)
+
+            if self.gradient_checkpointing and self.training:
+                layer_outputs = self._gradient_checkpointing_func(
+                    decoder_layer.__call__,
+                    hidden_states,
+                    attention_mask,
+                    position_ids,
+                    past_key_values,
+                    output_attentions,
+                    use_cache,
+                )
+            else:
+                layer_outputs = decoder_layer(
+                    hidden_states,
+                    attention_mask=attention_mask,
+                    position_ids=position_ids,
+                    past_key_value=past_key_values,
+                    output_attentions=output_attentions,
+                    use_cache=use_cache,
+                )
+
+            hidden_states = layer_outputs[0]
+
+            if use_cache:
+                next_decoder_cache = layer_outputs[2 if output_attentions else 1]
+
+            if output_attentions:
+                all_self_attns += (layer_outputs[1],)
+
+        hidden_states = self.norm(hidden_states)
+
+        # add hidden states from the last decoder layer
+        if output_hidden_states:
+            all_hidden_states += (hidden_states,)
+
+        next_cache = None
+        if use_cache:
+            next_cache = next_decoder_cache.to_legacy_cache() if use_legacy_cache else next_decoder_cache
+
+        if not return_dict:
+            return tuple(v for v in [hidden_states, next_cache, all_hidden_states, all_self_attns] if v is not None)
+        return BaseModelOutputWithPast(
+            last_hidden_state=hidden_states,
+            past_key_values=next_cache,
+            hidden_states=all_hidden_states,
+            attentions=all_self_attns,
+        )
+
+
+class Qwen2ForCausalLM(Qwen2PreTrainedModel):
+    _tied_weights_keys = ["lm_head.weight"]
+
+    def __init__(self, config):
+        super().__init__(config)
+        self.model = Qwen2Model(config)
+        self.vocab_size = config.vocab_size
+        self.lm_head = nn.Linear(config.hidden_size, config.vocab_size, bias=False)
+
+        self.text_mask_token_id = getattr(config, 'text_mask_token_id', 151676)
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+
+    def get_input_embeddings(self):
+        return self.model.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.model.embed_tokens = value
+
+    def get_output_embeddings(self):
+        return self.lm_head
+
+    def set_output_embeddings(self, new_embeddings):
+        self.lm_head = new_embeddings
+
+    def set_decoder(self, decoder):
+        self.model = decoder
+
+    def get_decoder(self):
+        return self.model
+
+    @add_start_docstrings_to_model_forward(QWEN2_INPUTS_DOCSTRING)
+    @replace_return_docstrings(output_type=CausalLMOutputWithPast, config_class=_CONFIG_FOR_DOC)
+    def forward(
+        self,
+        input_ids: torch.LongTensor = None,
+        visual_features: Optional[torch.FloatTensor] = None,
+        image_token_index: int = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[List[torch.FloatTensor]] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        labels: Optional[torch.LongTensor] = None,
+        use_cache: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        return_dict: Optional[bool] = None,
+    ) -> Union[Tuple, CausalLMOutputWithPast]:
+        r"""
+        Args:
+            labels (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+                Labels for computing the masked language modeling loss. Indices should either be in `[0, ...,
+                config.vocab_size]` or -100 (see `input_ids` docstring). Tokens with indices set to `-100` are ignored
+                (masked), the loss is only computed for the tokens with labels in `[0, ..., config.vocab_size]`.
+
+        Returns:
+
+        Example:
+
+        ```python
+        >>> from transformers import AutoTokenizer, Qwen2ForCausalLM
+
+        >>> model = Qwen2ForCausalLM.from_pretrained(PATH_TO_CONVERTED_WEIGHTS)
+        >>> tokenizer = AutoTokenizer.from_pretrained(PATH_TO_CONVERTED_TOKENIZER)
+
+        >>> prompt = "Hey, are you conscious? Can you talk to me?"
+        >>> inputs = tokenizer(prompt, return_tensors="pt")
+
+        >>> # Generate
+        >>> generate_ids = model.generate(inputs.input_ids, max_length=30)
+        >>> tokenizer.batch_decode(generate_ids, skip_special_tokens=True, clean_up_tokenization_spaces=False)[0]
+        "Hey, are you conscious? Can you talk to me?\nI'm not conscious, but I can talk to you."
+        ```"""
+
+        output_attentions = output_attentions if output_attentions is not None else self.config.output_attentions
+        output_hidden_states = (
+            output_hidden_states if output_hidden_states is not None else self.config.output_hidden_states
+        )
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        # decoder outputs consists of (dec_features, layer_state, dec_hidden, dec_attn)
+        outputs = self.model(
+            input_ids=input_ids,
+            visual_features=visual_features,
+            image_token_index=image_token_index,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_values=past_key_values,
+            inputs_embeds=inputs_embeds,
+            use_cache=use_cache,
+            output_attentions=output_attentions,
+            output_hidden_states=output_hidden_states,
+            return_dict=return_dict,
+        )
+
+        hidden_states = outputs[0]
+        logits = self.lm_head(hidden_states)
+        logits = logits.float()
+
+        loss = None
+        if labels is not None:
+
+            # Shift so that tokens < n predict n
+            shift_logits = logits[..., :-1, :].contiguous()
+            shift_labels = labels[..., 1:].contiguous()
+
+            # Flatten the tokens
+            loss_fct = CrossEntropyLoss()
+            shift_logits = shift_logits.view(-1, self.config.vocab_size)
+
+            shift_labels = shift_labels.view(-1)
+            shift_labels = shift_labels.to(shift_logits.device)
+            loss = loss_fct(shift_logits, shift_labels)
+
+            pos_masks = find_pred_pos_from_input_ids(input_ids, text_mask_token_id=self.text_mask_token_id)
+            shift_input_ids = input_ids[..., :-1].contiguous()
+            shift_pos_masks = pos_masks[:, :-1]
+            shift_input_ids = shift_input_ids.view(-1)
+            max_n_future_tokens = min(4, self.model.block_size)
+            pos_loss_list = torch.zeros(max_n_future_tokens, device=shift_logits.device)
+            shift_pos_masks = shift_pos_masks.reshape(-1)
+
+            for ix in range(max_n_future_tokens):
+                seg_loss = F.cross_entropy(
+                    shift_logits[shift_pos_masks == ix],
+                    shift_labels[shift_pos_masks == ix],
+                    reduction='mean'
+                )
+                pos_loss_list[ix] = seg_loss
+
+
+        if not return_dict:
+            output = (logits,) + outputs[1:]
+            return (loss,) + output if loss is not None else output
+        
+        if self.training:
+            return CausalLMOutputWithPast(
+                loss=loss,
+                logits=logits,
+                past_key_values=outputs.past_key_values,
+                hidden_states=outputs.hidden_states,
+                attentions=outputs.attentions,
+            ), pos_loss_list
+
+        return CausalLMOutputWithPast(
+            loss=loss,
+            logits=logits,
+            past_key_values=outputs.past_key_values,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )
+
+    def prepare_inputs_for_generation(
+        self, input_ids, past_key_values=None, attention_mask=None, inputs_embeds=None, **kwargs
+    ):
+        # Omit tokens covered by past_key_values
+        if past_key_values is not None:
+            if isinstance(past_key_values, Cache):
+                cache_length = past_key_values.get_seq_length()
+                past_length = past_key_values.seen_tokens
+                max_cache_length = past_key_values.get_max_length()
+            else:
+                cache_length = past_length = past_key_values[0][0].shape[2]
+                max_cache_length = None
+
+            # Keep only the unprocessed tokens:
+            # 1 - If the length of the attention_mask exceeds the length of input_ids, then we are in a setting where
+            # some of the inputs are exclusively passed as part of the cache (e.g. when passing input_embeds as
+            # input)
+            if attention_mask is not None and attention_mask.shape[1] > input_ids.shape[1]:
+                input_ids = input_ids[:, -(attention_mask.shape[1] - past_length) :]
+            # 2 - If the past_length is smaller than input_ids', then input_ids holds all input tokens. We can discard
+            # input_ids based on the past_length.
+            elif past_length < input_ids.shape[1]:
+                input_ids = input_ids[:, past_length:]
+            # 3 - Otherwise (past_length >= input_ids.shape[1]), let's assume input_ids only has unprocessed tokens.
+
+            # If we are about to go beyond the maximum cache length, we need to crop the input attention mask.
+            if (
+                max_cache_length is not None
+                and attention_mask is not None
+                and cache_length + input_ids.shape[1] > max_cache_length
+            ):
+                attention_mask = attention_mask[:, -max_cache_length:]
+
+        position_ids = kwargs.get("position_ids", None)
+        if attention_mask is not None and position_ids is None:
+            # create position_ids on the fly for batch generation
+            position_ids = attention_mask.long().cumsum(-1) - 1
+            position_ids.masked_fill_(attention_mask == 0, 1)
+            if past_key_values:
+                position_ids = position_ids[:, -input_ids.shape[1] :]
+
+        # if `inputs_embeds` are passed, we only want to use them in the 1st generation step
+        if inputs_embeds is not None and past_key_values is None:
+            model_inputs = {"inputs_embeds": inputs_embeds}
+        else:
+            model_inputs = {"input_ids": input_ids}
+
+        model_inputs.update(
+            {
+                "position_ids": position_ids,
+                "past_key_values": past_key_values,
+                "use_cache": kwargs.get("use_cache"),
+                "attention_mask": attention_mask,
+            }
+        )
+        return model_inputs
+
+    @staticmethod
+    def _reorder_cache(past_key_values, beam_idx):
+        reordered_past = ()
+        for layer_past in past_key_values:
+            reordered_past += (
+                tuple(past_state.index_select(0, beam_idx.to(past_state.device)) for past_state in layer_past),
+            )
+        return reordered_past
+
+
+@add_start_docstrings(
+    """
+    The Qwen2 Model transformer with a sequence classification head on top (linear layer).
+
+    [`Qwen2ForSequenceClassification`] uses the last token in order to do the classification, as other causal models
+    (e.g. GPT-2) do.
+
+    Since it does classification on the last token, it requires to know the position of the last token. If a
+    `pad_token_id` is defined in the configuration, it finds the last token that is not a padding token in each row. If
+    no `pad_token_id` is defined, it simply takes the last value in each row of the batch. Since it cannot guess the
+    padding tokens when `inputs_embeds` are passed instead of `input_ids`, it does the same (take the last value in
+    each row of the batch).
+    """,
+    QWEN2_START_DOCSTRING,
+)
+class Qwen2ForSequenceClassification(Qwen2PreTrainedModel):
+    def __init__(self, config):
+        super().__init__(config)
+        self.num_labels = config.num_labels
+        self.model = Qwen2Model(config)
+        self.score = nn.Linear(config.hidden_size, self.num_labels, bias=False)
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    def get_input_embeddings(self):
+        return self.model.embed_tokens
+
+    def set_input_embeddings(self, value):
+        self.model.embed_tokens = value
+
+    @add_start_docstrings_to_model_forward(QWEN2_INPUTS_DOCSTRING)
+    def forward(
+        self,
+        input_ids: torch.LongTensor = None,
+        attention_mask: Optional[torch.Tensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[List[torch.FloatTensor]] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        labels: Optional[torch.LongTensor] = None,
+        use_cache: Optional[bool] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        return_dict: Optional[bool] = None,
+    ) -> Union[Tuple, SequenceClassifierOutputWithPast]:
+        r"""
+        labels (`torch.LongTensor` of shape `(batch_size,)`, *optional*):
+            Labels for computing the sequence classification/regression loss. Indices should be in `[0, ...,
+            config.num_labels - 1]`. If `config.num_labels == 1` a regression loss is computed (Mean-Square loss), If
+            `config.num_labels > 1` a classification loss is computed (Cross-Entropy).
+        """
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        transformer_outputs = self.model(
+            input_ids,
+            attention_mask=attention_mask,
+            position_ids=position_ids,
+            past_key_values=past_key_values,
+            inputs_embeds=inputs_embeds,
+            use_cache=use_cache,
+            output_attentions=output_attentions,
+            output_hidden_states=output_hidden_states,
+            return_dict=return_dict,
+        )
+        hidden_states = transformer_outputs[0]
+        logits = self.score(hidden_states)
+
+        if input_ids is not None:
+            batch_size = input_ids.shape[0]
+        else:
+            batch_size = inputs_embeds.shape[0]
+
+        if self.config.pad_token_id is None and batch_size != 1:
+            raise ValueError("Cannot handle batch sizes > 1 if no padding token is defined.")
+        if self.config.pad_token_id is None:
+            sequence_lengths = -1
+        else:
+            if input_ids is not None:
+                # if no pad token found, use modulo instead of reverse indexing for ONNX compatibility
+                sequence_lengths = torch.eq(input_ids, self.config.pad_token_id).int().argmax(-1) - 1
+                sequence_lengths = sequence_lengths % input_ids.shape[-1]
+                sequence_lengths = sequence_lengths.to(logits.device)
+            else:
+                sequence_lengths = -1
+
+        pooled_logits = logits[torch.arange(batch_size, device=logits.device), sequence_lengths]
+
+        loss = None
+        if labels is not None:
+            labels = labels.to(logits.device)
+            if self.config.problem_type is None:
+                if self.num_labels == 1:
+                    self.config.problem_type = "regression"
+                elif self.num_labels > 1 and (labels.dtype == torch.long or labels.dtype == torch.int):
+                    self.config.problem_type = "single_label_classification"
+                else:
+                    self.config.problem_type = "multi_label_classification"
+
+            if self.config.problem_type == "regression":
+                loss_fct = MSELoss()
+                if self.num_labels == 1:
+                    loss = loss_fct(pooled_logits.squeeze(), labels.squeeze())
+                else:
+                    loss = loss_fct(pooled_logits, labels)
+            elif self.config.problem_type == "single_label_classification":
+                loss_fct = CrossEntropyLoss()
+                loss = loss_fct(pooled_logits.view(-1, self.num_labels), labels.view(-1))
+            elif self.config.problem_type == "multi_label_classification":
+                loss_fct = BCEWithLogitsLoss()
+                loss = loss_fct(pooled_logits, labels)
+        if not return_dict:
+            output = (pooled_logits,) + transformer_outputs[1:]
+            return ((loss,) + output) if loss is not None else output
+
+        return SequenceClassifierOutputWithPast(
+            loss=loss,
+            logits=pooled_logits,
+            past_key_values=transformer_outputs.past_key_values,
+            hidden_states=transformer_outputs.hidden_states,
+            attentions=transformer_outputs.attentions,
+        )

modeling_vit.py

@@ -0,0 +1,615 @@
+# Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
+#
+# NVIDIA CORPORATION and its licensors retain all intellectual property
+# and proprietary rights in and to this software, related documentation
+# and any modifications thereto.  Any use, reproduction, disclosure or
+# distribution of this software and related documentation without an express
+# license agreement from NVIDIA CORPORATION is strictly prohibited.
+
+import math
+from copy import deepcopy
+from typing import Union, Tuple, Sequence, Optional, List
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+try:
+    from transformers.activations import PytorchGELUTanh
+except ImportError:
+    PytorchGELUTanh = lambda: nn.GELU(approximate='tanh')
+from transformers.modeling_utils import PreTrainedModel
+from transformers.utils import is_flash_attn_2_available, logging
+
+if is_flash_attn_2_available():
+    from flash_attn import flash_attn_varlen_func
+else:
+    flash_attn_varlen_func = None
+
+from transformers.configuration_utils import PretrainedConfig
+
+
+logger = logging.get_logger(__name__)
+
+
+class MoonViTConfig(PretrainedConfig):
+    model_type = "moonvit"
+
+    def __init__(
+        self,
+        patch_size: int = 14,
+        init_pos_emb_height: int = 64,
+        init_pos_emb_width: int = 64,
+        num_attention_heads: int = 16,
+        num_hidden_layers: int = 27,
+        hidden_size: int = 1152,
+        intermediate_size: int = 4304,
+        merge_kernel_size: tuple[int, int] = (2, 2),
+        **kwargs,
+    ):
+        super().__init__(**kwargs)
+        self.patch_size = patch_size
+        # Positional embedding config
+        self.init_pos_emb_height = init_pos_emb_height
+        self.init_pos_emb_width = init_pos_emb_width
+        # Transformer config
+        self.num_hidden_layers = num_hidden_layers
+        self.num_attention_heads = num_attention_heads
+        self.hidden_size = hidden_size
+        self.intermediate_size = intermediate_size
+        # Patch merger config
+        self.merge_kernel_size = merge_kernel_size
+
+
+def multihead_attention(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q_cu_seqlens: Optional[torch.Tensor] = None,
+    k_cu_seqlens: Optional[torch.Tensor] = None,
+):
+    """Multi-head attention using flash attention 2.
+
+    Args:
+        q, k, v: tensor of shape (batch_size, seqlen, num_heads, head_dim),
+            or (tot_seqlens, num_heads, head_dim) if packing.
+        q_cu_seqlens (torch.Tensor): cumulative sequence lengths of q.
+            The first element should be 0 and the last element should be q.shape[0].
+        k_cu_seqlens (torch.Tensor): cumulative sequence lengths of k.
+            The first element should be 0 and the last element should be k.shape[0].
+
+    Returns:
+        output: shape (batch_size, seqlen, dim) or (tot_seqlens, dim) if packing,
+            where dim = num_heads * head_dim
+    """
+    if flash_attn_varlen_func is None:
+        logger.warning_once(
+            "flash_attn is not available for MoonViT; falling back to sdpa attention."
+        )
+        return sdpa_attention(
+            q,
+            k,
+            v,
+            q_cu_seqlens=q_cu_seqlens,
+            k_cu_seqlens=k_cu_seqlens,
+        )
+
+    # Unified format legal check
+    assert q.dim() == k.dim() == v.dim() == 3, "q, k, v must have 3 dims"
+    assert q_cu_seqlens[-1] == q.shape[0], "q_cu_seqlens must sum to q.shape[0]"
+    assert (
+        k_cu_seqlens[-1] == k.shape[0] == v.shape[0]
+    ), "k_cu_seqlens must sum to k.shape[0]"
+    assert q.dtype in [
+        torch.bfloat16,
+        torch.float16,
+    ], f"unsupported dtype {q.dtype} for multihead attn"
+
+    max_seqlen_q = (q_cu_seqlens[1:] - q_cu_seqlens[:-1]).max().item()
+    max_seqlen_k = (k_cu_seqlens[1:] - k_cu_seqlens[:-1]).max().item()
+    attn_out = flash_attn_varlen_func(
+        q,
+        k,
+        v,
+        q_cu_seqlens,
+        k_cu_seqlens,
+        max_seqlen_q,
+        max_seqlen_k,
+        causal=False,
+    )
+    attn_out = attn_out.flatten(start_dim=-2)
+
+    return attn_out
+
+
+def sdpa_attention(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q_cu_seqlens: Optional[torch.Tensor] = None,
+    k_cu_seqlens: Optional[torch.Tensor] = None,
+) -> torch.Tensor:
+    """SDPA attention.
+
+    Args:
+        q, k, v: tensor of shape (batch_size, seqlen, num_heads, head_dim),
+            or (tot_seqlens, num_heads, head_dim) if packing.
+    """
+    seq_length = q.shape[0]
+    attention_mask = torch.zeros(
+        [1, seq_length, seq_length], device=q.device, dtype=torch.bool
+    )
+    for i in range(1, len(q_cu_seqlens)):
+        attention_mask[
+            ...,
+            q_cu_seqlens[i - 1] : q_cu_seqlens[i],
+            q_cu_seqlens[i - 1] : q_cu_seqlens[i],
+        ] = True
+    q = q.transpose(0, 1)
+    k = k.transpose(0, 1)
+    v = v.transpose(0, 1)
+    attn_output = F.scaled_dot_product_attention(q, k, v, attention_mask, dropout_p=0.0)
+    attn_output = attn_output.transpose(0, 1)
+    attn_output = attn_output.reshape(seq_length, -1)
+    return attn_output
+
+
+def eager_attention(
+    q: torch.Tensor,
+    k: torch.Tensor,
+    v: torch.Tensor,
+    q_cu_seqlens: Optional[torch.Tensor] = None,
+    k_cu_seqlens: Optional[torch.Tensor] = None,
+) -> torch.Tensor:
+    seq_length = q.shape[0]
+    attention_mask = torch.zeros(
+        [1, seq_length, seq_length], device=q.device, dtype=torch.bool
+    )
+    for i in range(1, len(q_cu_seqlens)):
+        attention_mask[
+            ...,
+            q_cu_seqlens[i - 1] : q_cu_seqlens[i],
+            q_cu_seqlens[i - 1] : q_cu_seqlens[i],
+        ] = True
+    q = q.transpose(0, 1)
+    k = k.transpose(0, 1)
+    v = v.transpose(0, 1)
+
+    attn_weight = q @ k.transpose(-2, -1) / math.sqrt(q.shape[-1])
+    attn_weight += attention_mask
+    attn_weight = torch.softmax(attn_weight, dim=-1, dtype=torch.float32).to(q.dtype)
+
+    attn_output = attn_weight @ v
+    attn_output = attn_output.transpose(0, 1)
+    attn_output = attn_output.reshape(seq_length, -1)
+    return attn_output
+
+
+VL_VISION_ATTENTION_FUNCTIONS = {
+    "flash_attention_2": multihead_attention,
+    "sdpa": sdpa_attention,
+    "eager": eager_attention,
+}
+
+
+def _apply_rope_input_validation(x, freqs_cis):
+    assert x.ndim == freqs_cis.ndim + 1, (x.shape, freqs_cis.shape)
+    assert x.shape[:-2] == freqs_cis.shape[:-1], (x.shape, freqs_cis.shape)
+    assert x.shape[-1] == 2 * freqs_cis.shape[-1], (x.shape, freqs_cis.shape)
+    assert freqs_cis.dtype == torch.complex64, freqs_cis.dtype
+
+
+def apply_rope(
+    xq: torch.Tensor, xk: torch.Tensor, freqs_cis: torch.Tensor
+) -> tuple[torch.Tensor, torch.Tensor]:
+    """
+    Args: (The leading dimensions of all inputs should be the same)
+        xq: query, tensor of shape (..., num_heads, head_dim)
+        xk: key, tensor of shape (..., num_heads, head_dim)
+        freqs_cis: tensor of shape (..., head_dim/2), dtype=torch.complex64. It contains the precomputed cis(freqs) for each position in the 2D grid.
+    Returns:
+        xq_out, xk_out: tensors of shape (..., num_heads, head_dim)
+    """
+    _apply_rope_input_validation(xq, freqs_cis)
+    _apply_rope_input_validation(xk, freqs_cis)
+
+    freqs_cis = freqs_cis.unsqueeze(-2)  # ..., 1, head_dim/2
+    # ..., num_heads, head_dim/2
+    xq_ = torch.view_as_complex(xq.float().view(*xq.shape[:-1], -1, 2))
+    xk_ = torch.view_as_complex(xk.float().view(*xq.shape[:-1], -1, 2))
+    xq_out = torch.view_as_real(xq_ * freqs_cis).flatten(-2)  # ..., num_heads, head_dim
+    xk_out = torch.view_as_real(xk_ * freqs_cis).flatten(-2)  # ..., num_heads, head_dim
+    return xq_out.type_as(xq), xk_out.type_as(xk)
+
+
+class Learnable2DInterpPosEmb(nn.Module):
+    def __init__(
+        self, height: int, width: int, dim: int, interpolation_mode: str = "bicubic"
+    ) -> None:
+        super().__init__()
+        self.height = height
+        self.width = width
+        self.interpolation_mode = interpolation_mode
+        self.weight = nn.Parameter(torch.empty(height, width, dim))
+        self.reset_parameters()
+
+    def reset_parameters(self):
+        nn.init.normal_(self.weight)
+
+    def forward(self, x: torch.Tensor, grid_hws: torch.Tensor) -> torch.Tensor:
+        pos_embs = []
+        for shape in grid_hws.tolist():
+            if shape == self.weight.shape[:-1]:
+                pos_embs.append(self.weight.flatten(end_dim=1))
+            else:
+                pos_embs.append(
+                    F.interpolate(
+                        self.weight.permute((2, 0, 1)).unsqueeze(0),
+                        size=shape,
+                        mode=self.interpolation_mode,
+                    )
+                    .squeeze(0)
+                    .permute((1, 2, 0))
+                    .flatten(end_dim=1)
+                )
+        out = x + torch.cat(pos_embs)
+        return out
+
+
+class MoonVisionPatchEmbed(nn.Module):
+
+    def __init__(
+        self,
+        out_dim: int,
+        in_dim: int = 3,
+        patch_size: Union[int, Tuple[int, int]] = (14, 14),
+        pos_emb_height: int = 14,
+        pos_emb_width: int = 14,
+    ):
+        super().__init__()
+        assert isinstance(
+            patch_size, (int, Sequence)
+        ), f"Invalid patch_size type: {type(patch_size)}"
+        if isinstance(patch_size, int):
+            patch_size = (patch_size, patch_size)
+        assert (
+            len(patch_size) == 2
+        ), f"Expected patch_size to be a tuple of 2, got {patch_size}"
+        self.patch_size = patch_size
+
+        self.proj = nn.Conv2d(
+            in_dim, out_dim, kernel_size=patch_size, stride=patch_size
+        )
+
+        self.pos_emb = Learnable2DInterpPosEmb(
+            height=pos_emb_height, width=pos_emb_width, dim=out_dim
+        )
+
+    def forward(self, x: torch.Tensor, grid_hws: torch.Tensor) -> torch.Tensor:
+        """
+        Args:
+            x (L, Channels): input tensor
+            grid_hws (N, 2): grid height and width
+
+        Returns:
+            (L, Cout) tensor
+        """
+        x = self.proj(x).view(x.size(0), -1)
+        # apply positional embedding
+        x = self.pos_emb(x, grid_hws)
+        return x
+
+
+class Rope2DPosEmb(nn.Module):
+    """2D rotary position embedding with multi-resolution support.
+
+    This class is intended to be used in the following way:
+    1. Before training, create an instance of Rope2DPosEmb. This instance will hold the precomputed cis.
+    2. Before each forward pass, call `get_freqs_cis_by_*` to get the `freqs_cis` tensor for this iteration.
+    3. During the forward pass, pass the `freqs_cis` tensor to each attention layer, and call `apply` just before each attention operation.
+        The rope is shared across all attention layers and all heads.
+
+    Refs:
+    - RoFormer: https://arxiv.org/abs/2104.09864
+    - VisionLLaMA: https://arxiv.org/abs/2403.00522
+    - https://github.com/Meituan-AutoML/VisionLLaMA/blob/main/dit/models.py
+
+    Args:
+        dim (int): usually the multi-head attention dimension, should be divisible by 4 (TODO: relax this constraint if needed)
+        max_height (int): the maximum height of the 2D grid
+        max_width (int): the maximum width of the 2D grid
+        theta_base (float): the base of the theta
+        device (str): the device to store the precomputed cis
+    """
+
+    def __init__(self, dim: int, max_height: int, max_width: int, theta_base=10000):
+        super().__init__()
+        self.dim = dim
+        assert self.dim % 4 == 0, "dim must be divisible by 4"
+        self.max_height = max_height
+        self.max_width = max_width
+        self.theta_base = theta_base
+
+        self.freqs_cis = None
+
+    def extra_repr(self):
+        return f"dim={self.dim}, max_height={self.max_height}, max_width={self.max_width}, theta_base={self.theta_base}"
+
+    def _precompute_freqs_cis(self, device: torch.device) -> torch.Tensor:
+        """Calculate the cis(freqs) for each position in the 2D grid.
+
+        Return: complex tensor of shape (max_height, max_width, dim//2) and value:
+            height axis: ret[h, w, 2*i] = cis(h * theta_base**(-4*i/dim))
+            weight axis: ret[h, w, 2*i+1] = cis(w * theta_base**(-4*i/dim))   with (i in [0, dim//4))
+            note: `cis` is a mathematical notation defined by cis x = cos x + i sin x,
+        """
+        N = self.max_height * self.max_width
+        flat_pos = torch.arange(0, N).float().to(device)
+        x_pos = flat_pos % self.max_width
+        y_pos = flat_pos // self.max_width
+        dim_range = (
+            torch.arange(0, self.dim, 4)[: (self.dim // 4)].float().to(device)
+        )  # C/4
+        freqs = 1.0 / (self.theta_base ** (dim_range / self.dim))
+        x_freqs = torch.outer(x_pos, freqs).float()  # N, C/4
+        y_freqs = torch.outer(y_pos, freqs).float()  # N, C/4
+        x_cis = torch.polar(torch.ones_like(x_freqs), x_freqs)  # N, C/4
+        y_cis = torch.polar(torch.ones_like(y_freqs), y_freqs)  # N, C/4
+        # N, C/4, 2
+        freqs_cis = torch.cat(
+            [x_cis.unsqueeze(dim=-1), y_cis.unsqueeze(dim=-1)], dim=-1
+        )
+        # max_height, max_width, C/2
+        freqs_cis = freqs_cis.reshape(self.max_height, self.max_width, -1)
+        return freqs_cis
+
+    def get_freqs_cis(self, grid_hws: torch.Tensor) -> torch.Tensor:
+        """
+        Args:
+            grid_hws (torch.Tensor): grid height and width
+
+        Returns:
+            freqs_cis: tensor of shape (sum(t * height * width), dim//2)
+        """
+        if self.freqs_cis is None:
+            self.freqs_cis = self._precompute_freqs_cis(grid_hws.device)
+
+        shapes = grid_hws.tolist()
+        assert all(
+            1 <= h <= self.max_height and 1 <= w <= self.max_width for h, w in shapes
+        ), (
+            shapes,
+            self.max_height,
+            self.max_width,
+        )
+        freqs_cis = torch.cat(
+            [self.freqs_cis[:h, :w].reshape(-1, self.dim // 2) for h, w in shapes],
+            dim=0,
+        )
+        return freqs_cis
+
+
+class MLP2(nn.Module):
+    """
+    Args:
+        dims: [in_dim, hidden_dim, out_dim]
+        bias: whether to use bias in linear layer.
+    """
+
+    def __init__(self, dims: list[int], activation, bias=True):
+        super().__init__()
+        assert len(dims) == 3
+        self.fc0 = nn.Linear(dims[0], dims[1], bias=bias)
+        self.fc1 = nn.Linear(dims[1], dims[2], bias=bias)
+        self.activation = activation
+        for m in [self.fc0, self.fc1]:
+            nn.init.trunc_normal_(m.weight, std=math.sqrt(2 / m.in_features))
+            if m.bias is not None:
+                nn.init.zeros_(m.bias)
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        x = self.fc0(x)
+        x = self.activation(x)
+        return self.fc1(x)
+
+
+class MoonVitEncoderLayer(nn.Module):
+
+    def __init__(
+        self,
+        num_heads: int,
+        hidden_dim: int,
+        mlp_dim: int,
+        *,
+        attn_implementation: str = "eager",
+        activation=F.gelu,
+        attn_bias: bool = False,
+    ):
+        super().__init__()
+        self.num_heads = num_heads
+        self.hidden_dim = hidden_dim
+        self.hidden_size_per_attention_head = self.hidden_dim // self.num_heads
+        self.attn_implementation = attn_implementation
+
+        self.norm0 = nn.LayerNorm(hidden_dim)
+        self.norm1 = nn.LayerNorm(hidden_dim)
+        self.mlp = MLP2([hidden_dim, mlp_dim, hidden_dim], activation)
+        self.wqkv = nn.Linear(hidden_dim, hidden_dim * 3, bias=attn_bias)
+        self.wo = nn.Linear(hidden_dim, hidden_dim, bias=attn_bias)
+
+    def attention_qkvpacked(
+        self,
+        x: torch.Tensor,
+        cu_seqlens: torch.Tensor,
+        rope_freqs_cis: Optional[torch.Tensor] = None,
+    ):
+        """
+        Args:
+            x (torch.Tensor): (batch_size, seqlen, hidden_dim)
+            cu_seqlens (torch.Tensor):
+        """
+        xqkv = self.wqkv(x)
+
+        qkv_shape = xqkv.size()[:-1] + (
+            3,
+            self.num_heads,
+            self.hidden_size_per_attention_head,
+        )
+        # xqkv: (batch_size, seqlen, 3, nheads, headdim)
+        xqkv = xqkv.view(*qkv_shape)
+        xq, xk, xv = torch.unbind(xqkv, dim=-3)
+
+        xq, xk = apply_rope(xq, xk, rope_freqs_cis)
+
+        attn_func = VL_VISION_ATTENTION_FUNCTIONS[self.attn_implementation]
+        attn_out = attn_func(
+            xq, xk, xv, q_cu_seqlens=cu_seqlens, k_cu_seqlens=cu_seqlens
+        )
+
+        attn_out = self.wo(attn_out)
+        return attn_out
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        cu_seqlens: torch.Tensor,
+        rope_freqs_cis: Union[torch.Tensor, None] = None,
+    ) -> torch.Tensor:
+        """
+        Args:
+            hidden_states: non-packed (B, N, D) or packed (L, D). if non-packed, seqlens should be None, if packed, seqlens should be set
+
+        Returns:
+            output: same shape of input, non-packed (B, N, D) for non-packed input, (L, D) for packed input
+        """
+        residual = hidden_states
+        hidden_states = self.norm0(hidden_states)
+        attn_out = self.attention_qkvpacked(
+            hidden_states, cu_seqlens, rope_freqs_cis=rope_freqs_cis
+        )
+        hidden_states = residual + attn_out
+
+        residual = hidden_states
+        hidden_states = self.mlp(self.norm1(hidden_states))
+        hidden_states = residual + hidden_states
+        return hidden_states
+
+
+class MoonVitEncoder(nn.Module):
+
+    def __init__(
+        self,
+        hidden_dim: int,
+        num_layers: int,
+        block_cfg: dict,
+    ) -> None:
+        super().__init__()
+
+        self.rope_2d = Rope2DPosEmb(
+            block_cfg["hidden_dim"] // block_cfg["num_heads"], 512, 512
+        )
+        self.blocks = nn.ModuleList(
+            [MoonVitEncoderLayer(**block_cfg) for _ in range(num_layers)]
+        )
+        self.final_layernorm = nn.LayerNorm(hidden_dim)
+
+    def forward(
+        self, hidden_states: torch.Tensor, grid_hws: torch.Tensor
+    ) -> torch.Tensor:
+        rope_freqs_cis = self.rope_2d.get_freqs_cis(grid_hws=grid_hws)
+
+        lengths = torch.cat(
+            (
+                torch.zeros(1, device=hidden_states.device, dtype=grid_hws.dtype),
+                grid_hws[:, 0] * grid_hws[:, 1],
+            )
+        )
+        cu_seqlens = lengths.cumsum(dim=0, dtype=torch.int32)
+
+        for _, block in enumerate(self.blocks):
+            hidden_states = block(
+                hidden_states, cu_seqlens, rope_freqs_cis=rope_freqs_cis
+            )
+
+        hidden_states = self.final_layernorm(hidden_states)
+
+        return hidden_states
+
+
+def patch_merger(
+    x: torch.Tensor,
+    grid_hws: torch.Tensor,
+    merge_kernel_size: list[int, int] = (2, 2),
+) -> List[torch.Tensor]:
+    d_model = x.size(-1)
+
+    outputs = []
+    pre_sum = 0
+    for x_shape in grid_hws.tolist():
+        height, width = x_shape[0], x_shape[1]
+        # Get the current sequence
+        seq = x[pre_sum : pre_sum + height * width]
+        # Reshape along self.merge_kernel_size and concat to the last dimension
+        kernel_height, kernel_width = merge_kernel_size
+        new_height, new_width = height // kernel_height, width // kernel_width
+        reshaped_seq = seq.view(
+            new_height, kernel_height, new_width, kernel_width, d_model
+        )
+        reshaped_seq = reshaped_seq.permute(0, 2, 1, 3, 4).contiguous()
+        padded_seq = reshaped_seq.view(
+            new_height * new_width, -1
+        )
+        outputs.append(padded_seq)
+        pre_sum += height * width
+
+    return outputs
+
+
+class MoonVitPretrainedModel(PreTrainedModel):
+    config_class = MoonViTConfig
+    model_type = "moonvit"
+    _no_split_modules = ["PackingTransformer"]
+    _supports_flash_attn_2 = True
+    _supports_sdpa = True
+
+    def __init__(self, config: MoonViTConfig, *inputs, **kwargs):
+        super().__init__(config, *inputs, **kwargs)
+        config = deepcopy(config)
+        self.merge_kernel_size = config.merge_kernel_size
+        self.patch_size = config.patch_size
+        self.patch_embed = MoonVisionPatchEmbed(
+            out_dim=config.hidden_size,
+            patch_size=config.patch_size,
+            pos_emb_height=config.init_pos_emb_height,
+            pos_emb_width=config.init_pos_emb_width,
+        )
+
+        self.encoder = MoonVitEncoder(
+            hidden_dim=config.hidden_size,
+            num_layers=config.num_hidden_layers,
+            block_cfg={
+                "num_heads": config.num_attention_heads,
+                "hidden_dim": config.hidden_size,
+                "mlp_dim": config.intermediate_size,
+                "activation": PytorchGELUTanh(),
+                "attn_bias": True,
+                "attn_implementation": config._attn_implementation,
+            },
+        )
+
+    def forward(
+        self, pixel_values: torch.Tensor, grid_hws: torch.Tensor
+    ) -> torch.Tensor:
+        """
+        Args:
+            pixel_values (torch.Tensor): The input pixel values.
+            grid_hws (torch.Tensor): The grid height and width.
+
+        Returns:
+            torch.Tensor: The output tokens.
+        """
+        hidden_states = self.patch_embed(pixel_values, grid_hws)
+        hidden_states = self.encoder(hidden_states, grid_hws)
+        hidden_states = patch_merger(
+            hidden_states, grid_hws, merge_kernel_size=self.merge_kernel_size
+        )
+        return hidden_states

preprocessor_config.json

@@ -0,0 +1,24 @@
+{
+  "auto_map": {
+    "AutoImageProcessor": "image_processing_locateanything.LocateAnythingImageProcessor",
+    "AutoProcessor": "processing_locateanything.LocateAnythingProcessor"
+  },
+  "image_mean": [
+    0.5,
+    0.5,
+    0.5
+  ],
+  "image_processor_type": "LocateAnythingImageProcessor",
+  "image_std": [
+    0.5,
+    0.5,
+    0.5
+  ],
+  "in_token_limit": 25600,
+  "merge_kernel_size": [
+    2,
+    2
+  ],
+  "patch_size": 14,
+  "processor_class": "LocateAnythingProcessor"
+}

processing_locateanything.py

@@ -0,0 +1,678 @@
+# coding=utf-8
+# Copyright 2024 The HuggingFace Inc. team.
+#
+# 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.
+"""
+Processor class for LocateAnything.
+"""
+
+import math
+import os
+from typing import Iterable, List, Union, Literal
+import base64
+import sys
+import time
+import warnings
+from functools import lru_cache
+from io import BytesIO
+import re
+import requests
+import torch
+import torchvision
+from packaging import version
+from PIL import Image
+from torchvision import io
+from torchvision import transforms
+from torchvision.transforms import InterpolationMode
+from typing import Optional, Any
+import numpy as np
+
+from transformers.feature_extraction_utils import BatchFeature
+from transformers.image_utils import ImageInput
+try:
+    from transformers.image_utils import VideoInput
+except ImportError:
+    VideoInput = None
+from transformers.processing_utils import ProcessingKwargs, ProcessorMixin, Unpack
+from transformers.tokenization_utils_base import PreTokenizedInput, TextInput
+from transformers.utils import logging
+import lmdb
+import cv2
+import pickle
+import decord
+
+logger = logging.get_logger(__name__)
+
+FPS = 2.0
+MAX_FRAMES = 64 
+VIDEO_TOTAL_PIXELS = int(float(os.environ.get('VIDEO_MAX_PIXELS', 32000 * 28 * 28 * 0.9)))
+logger.info(f"set VIDEO_TOTAL_PIXELS: {VIDEO_TOTAL_PIXELS}")
+
+
+def to_rgb(pil_image: Image.Image) -> Image.Image:
+      if pil_image.mode == 'RGBA':
+          white_background = Image.new("RGB", pil_image.size, (255, 255, 255))
+          white_background.paste(pil_image, mask=pil_image.split()[3])  # Use alpha channel as mask
+          return white_background
+      else:
+          return pil_image.convert("RGB")
+
+def read_img_from_lmdb_v2(image_data):
+    # special case for AgiBotWorld
+    lmdb_file, lmdb_key = image_data['lmdb_file'], image_data['lmdb_key']
+    key = lmdb_key.encode('ascii')
+    env = lmdb.open(lmdb_file, max_readers=10240, readonly=True, lock=False, readahead=False, meminit=False)
+    txn = env.begin()
+    value = txn.get(key)
+    if value is None:
+        print(f"Warning: Key {key} not found.")
+        return None
+    record = pickle.loads(value)
+    image_bgr = cv2.imdecode(np.frombuffer(record['image'], dtype=np.uint8), cv2.IMREAD_COLOR)
+    image_rgb = cv2.cvtColor(image_bgr, cv2.COLOR_BGR2RGB)
+    image = Image.fromarray(image_rgb)
+        
+    return image
+
+def parse_lmdb_image_data(image_data):
+    lmdb_file = image_data['lmdb_file']
+    if not os.path.exists(lmdb_file):
+        if "/home/zhidingy/workspace/libs/eagle/Eagle2/" in lmdb_file:
+            image_data['lmdb_file'] = lmdb_file.replace("/home/zhidingy/workspace/libs/eagle/Eagle2/", "")
+        else:
+            raise ValueError(f"LMDB file {lmdb_file} does not exist")
+    # special case for AgiBotWorld
+    if 'AgiBotWorld' in image_data['lmdb_file']:
+        return read_img_from_lmdb_v2(image_data)
+    
+    try:
+        env = lmdb.open(image_data['lmdb_file'], readonly=True, lock=False, max_readers=10240)
+    except Exception as e:
+        print(f"Failed to open lmdb file {image_data['lmdb_file']}. Error message: {e}", flush=True)
+        raise e
+
+    with env.begin(write=False) as txn:
+        try:
+            image_bin = txn.get(image_data['lmdb_key'].encode('ascii'))
+            buf = BytesIO(image_bin)
+        except Exception as e:
+            print(f"Failed to get image from lmdb file {image_data['lmdb_file']}. Error message: {e}", flush=True)
+            raise e
+    try:
+        image = Image.open(buf)
+    except Exception as e:
+        image_np = np.frombuffer(image_bin, dtype=np.uint8)
+        image_bgr = cv2.imdecode(image_np, cv2.IMREAD_COLOR)
+        image_rgb = cv2.cvtColor(image_bgr, cv2.COLOR_BGR2RGB)
+        image = Image.fromarray(image_rgb)
+    return image
+
+def fetch_image(ele: dict[str, str | Image.Image]) -> Image.Image:
+    if "image" in ele:
+        image = ele["image"]
+    else:
+        image = ele["image_url"]
+    image_obj = None
+    if isinstance(image, Image.Image):
+        image_obj = image
+    elif isinstance(image, dict) and 'lmdb_file' in image:
+        image_obj = parse_lmdb_image_data(image)
+    elif image.startswith("http://") or image.startswith("https://"):
+        response = requests.get(image, stream=True)
+        image_obj = Image.open(BytesIO(response.content))
+    elif image.startswith("file://"):
+        image_obj = Image.open(image[7:])
+    elif image.startswith("data:image"):
+        if "base64," in image:
+            _, base64_data = image.split("base64,", 1)
+            data = base64.b64decode(base64_data)
+            image_obj = Image.open(BytesIO(data))
+    else:
+        image_obj = Image.open(image)
+    if image_obj is None:
+        raise ValueError(f"Unrecognized image input, support local path, http url, base64 and PIL.Image, got {image}")
+    image = to_rgb(image_obj)
+
+    return image
+
+
+def get_video_frame_indices(
+    ele: dict,
+    total_frames: int,
+    video_fps: int | float,
+) -> tuple[torch.Tensor, float]:
+    target_fps = ele.get("fps", FPS)
+    max_frames = ele.get("max_frames", MAX_FRAMES)
+    
+    nframes = (total_frames / video_fps) * target_fps
+    nframes = int(round(nframes))
+    nframes = max(1, nframes)
+    
+    if nframes > max_frames:
+        nframes = max_frames
+    
+    nframes = min(nframes, total_frames)
+    
+    if nframes == total_frames:
+        idx = torch.arange(total_frames).long()
+    else:
+        idx = torch.linspace(0, total_frames - 1, nframes).round().long()
+        
+    sample_fps = nframes / max(total_frames, 1e-6) * video_fps
+    
+    return idx, sample_fps
+
+def _read_video_torchvision(
+    ele: dict,
+) -> (torch.Tensor, float, list):
+    """read video using torchvision.io.read_video and return also per-frame timestamps"""
+    video_path = ele["video"]
+    if version.parse(torchvision.__version__) < version.parse("0.19.0"):
+        if "http://" in video_path or "https://" in video_path:
+            warnings.warn("torchvision < 0.19.0 does not support http/https video path, please upgrade to 0.19.0.")
+        if "file://" in video_path:
+            video_path = video_path[7:]
+    st = time.time()
+    
+    video, audio, info = io.read_video(
+        video_path,
+        start_pts=ele.get("video_start", 0.0),
+        end_pts=ele.get("video_end", None),
+        pts_unit="sec",
+        output_format="TCHW",
+    )
+    total_frames, video_fps = video.size(0), info["video_fps"]
+    logger.info(f"torchvision:  {video_path=}, {total_frames=}, {video_fps=}, time={time.time() - st:.3f}s")
+    
+    idx, sample_fps = get_video_frame_indices(ele, total_frames, video_fps)
+    
+    start_time = ele.get("video_start", 0.0)
+    timestamps = (start_time + idx.to(torch.float32) / video_fps).tolist()
+    
+    video = video[idx]
+    return video, sample_fps, timestamps
+
+
+def is_decord_available() -> bool:
+    import importlib.util
+    return importlib.util.find_spec("decord") is not None
+
+def _read_video_decord(
+    ele: dict,
+) -> (torch.Tensor, float, list):
+    """read video using decord.VideoReader and return also per-frame timestamps"""
+    video_path = ele["video"]
+    st = time.time()
+    vr = decord.VideoReader(video_path)
+    
+    total_frames, video_fps = len(vr), vr.get_avg_fps()
+    logger.info(f"decord:  {video_path=}, {total_frames=}, {video_fps=}, time={time.time() - st:.3f}s")
+    
+    idx_tensor, sample_fps = get_video_frame_indices(ele, total_frames, video_fps)
+    idx = idx_tensor.tolist()
+    
+    start_time = ele.get("video_start", 0.0)
+    timestamps = [start_time + i / video_fps for i in idx]
+    
+    video = vr.get_batch(idx).asnumpy()
+    video = torch.tensor(video).permute(0, 3, 1, 2)  # Convert to TCHW format
+    
+    return video, sample_fps, timestamps
+
+
+VIDEO_READER_BACKENDS = {
+    "decord": _read_video_decord,
+    "torchvision": _read_video_torchvision,
+}
+
+
+@lru_cache(maxsize=1)
+def get_video_reader_backend() -> str:
+    if is_decord_available():
+        video_reader_backend = "decord"
+    else:
+        video_reader_backend = "torchvision"
+    return video_reader_backend
+
+
+def fetch_video(ele: dict, return_video_sample_fps: bool = False, video_reader_backend: str = "torchvision") -> torch.Tensor | list[Image.Image]:
+    """
+    Fetches video, samples frames, resizes based on video_total_pixels, and returns as Tensor (TCHW).
+    """
+    if isinstance(ele["video"], str):
+        video_reader_backend = video_reader_backend if video_reader_backend is not None else get_video_reader_backend()
+        try:
+            video, sample_fps, timestamps = VIDEO_READER_BACKENDS[video_reader_backend](ele)
+        except Exception as e:
+            logger.warning(f"video_reader_backend {video_reader_backend} error, use torchvision as default, msg: {e}")
+            video, sample_fps, timestamps = VIDEO_READER_BACKENDS["torchvision"](ele)
+
+        nframes, _, height, width = video.shape
+        
+        video_total_pixels = ele.get("video_total_pixels", VIDEO_TOTAL_PIXELS)
+        current_pixels = nframes * height * width
+        
+        if current_pixels > video_total_pixels:
+            scale_factor = math.sqrt(video_total_pixels / current_pixels)
+            new_height = int(height * scale_factor)
+            new_width = int(width * scale_factor)
+            
+            video = transforms.functional.resize(
+                video,
+                [new_height, new_width],
+                interpolation=InterpolationMode.BICUBIC,
+                antialias=True,
+            ).float()
+        else:
+            video = video.float()
+
+        if return_video_sample_fps:
+            return video, sample_fps, timestamps
+        return video
+
+    else:
+        assert isinstance(ele["video"], (list, tuple))
+        process_info = ele.copy()
+        process_info.pop("type", None)
+        process_info.pop("video", None)
+        
+        images = [
+            fetch_image({"image": video_element, **process_info})
+            for video_element in ele["video"]
+        ]
+        
+        nframes = len(images)
+        timestamps = [-1 for i in range(nframes)] 
+        
+        # For list of images, we return list of PIL images directly, 
+        # the processor will handle conversion to tensor later.
+        if return_video_sample_fps:
+            return images, process_info.get("fps", 2.0), timestamps
+        return images
+
+class LocateAnythingProcessorKwargs(ProcessingKwargs, total=False):
+    _defaults = {
+        "text_kwargs": {
+            "padding": False,
+        },
+        "images_kwargs": {},
+        "videos_kwargs": {},
+    }
+
+
+class LocateAnythingProcessor(ProcessorMixin):
+    attributes = ["image_processor", "tokenizer"]
+    valid_kwargs = [
+        "chat_template",
+        "num_image_tokens",
+        "image_token",
+        "video_token",
+        "images_kwargs",
+        "videos_kwargs",
+        "text_kwargs",
+    ]
+    image_processor_class = "AutoImageProcessor" 
+    tokenizer_class = "AutoTokenizer"
+
+    def __init__(
+        self,
+        image_processor=None,
+        tokenizer=None,
+        chat_template=None,
+        image_token='<IMG_CONTEXT>',
+        video_token='<IMG_CONTEXT>',
+        merge_kernel_size=[2, 2], # Note: This might need adjustment based on your patch_size (14*14)
+        image_placeholder='image',
+        video_placeholder='video',
+        image_start_token='<img>',
+        image_end_token='</img>',
+        **kwargs,
+    ):  
+        self.image_token = tokenizer.image_token if hasattr(tokenizer, "image_token") else image_token
+        self.video_token = tokenizer.video_token if hasattr(tokenizer, "video_token") else video_token
+        self.image_token_id = (
+            tokenizer.image_token_id
+            if getattr(tokenizer, "image_token_id", None)
+            else tokenizer.convert_tokens_to_ids(self.image_token)
+        )
+        self.video_token_id = (
+            tokenizer.video_token_id
+            if getattr(tokenizer, "video_token_id", None)
+            else tokenizer.convert_tokens_to_ids(self.video_token)
+        )
+        self.image_placeholder = image_placeholder
+        self.video_placeholder = video_placeholder
+        self.merge_kernel_size = merge_kernel_size
+        self.image_start_token = image_start_token
+        self.image_end_token = image_end_token
+        if 'auto_map' in kwargs:
+            self.auto_map = kwargs['auto_map']
+        super().__init__(image_processor, tokenizer, chat_template=chat_template)
+
+    
+    def replace_media_placeholder(self, text, image_list, video_list, timestamps_list, fps_list, **output_kwargs):
+
+        num_of_images_in_this_sample = 0
+        num_of_videos_in_this_sample = 0
+        pattern = re.compile(rf"<({self.image_placeholder}|{self.video_placeholder})-(\d+)>")
+        unified_frame_list = []
+        
+        def replace_in_text(text):
+            def repl(match):
+                nonlocal unified_frame_list
+                nonlocal num_of_images_in_this_sample
+                nonlocal num_of_videos_in_this_sample
+                media_type = match.group(1)
+                idx_in_list = int(match.group(2)) - 1
+                idx_mapper = {0: "first", 1: "second", 2: "third", 3: "fourth", 4: "fifth", 5: "sixth", 6: "seventh", 7: "eighth", 8: "ninth", 9: "tenth"}  
+                
+                if media_type == 'image':
+                    # Call LocateAnythingImageProcessor with a single image in a list
+                    image_inputs = self.image_processor(images=[image_list[idx_in_list]], **output_kwargs["images_kwargs"])
+                    
+                    num_of_tokens_list = [int(h * w) // (self.image_processor.merge_kernel_size[0] * self.image_processor.merge_kernel_size[1]) for h, w in image_inputs['image_grid_hws']]
+                    
+                    special_placeholder = f"<image {idx_in_list+1}>{self.image_start_token}{self.image_token * num_of_tokens_list[0]}{self.image_end_token}"
+                    unified_frame_list.append(image_inputs)
+                    num_of_images_in_this_sample += 1
+                    
+                elif media_type == 'video':
+                    video_obj = video_list[idx_in_list]
+                    
+                    # Convert Tensor TCHW to list of PIL Images for the ImageProcessor
+                    if isinstance(video_obj, torch.Tensor):
+                        # video_obj is [T, C, H, W], float, likely 0-255 or standardized
+                        # LocateAnythingImageProcessor expects PIL or 0-255 inputs usually.
+                        # We need to convert back to PIL or List[Tensor] compatible with make_list_of_images
+                        video_frames = []
+                        for i in range(video_obj.shape[0]):
+                            frame = video_obj[i] # [C, H, W]
+                            # Assuming fetch_video returns float tensors.
+                            # If they are 0-255, convert to uint8.
+                            if frame.dtype.is_floating_point and frame.max() > 1.0:
+                                frame = frame.byte()
+                            elif frame.dtype.is_floating_point:
+                                frame = (frame * 255).byte()
+                            
+                            img = transforms.ToPILImage()(frame)
+                            video_frames.append(img)
+                    elif isinstance(video_obj, list):
+                        # Already list of PIL images
+                        video_frames = video_obj
+                    else:
+                        raise ValueError("Unsupported video format")
+
+                    # Call ImageProcessor with list of frames
+                    video_inputs = self.image_processor(images=video_frames, **output_kwargs["videos_kwargs"])
+                    
+                    # Calculate tokens per frame
+                    num_of_tokens_list = [int(h * w) // (self.image_processor.merge_kernel_size[0] * self.image_processor.merge_kernel_size[1]) for h, w in video_inputs['image_grid_hws']]
+                    
+                    if timestamps_list is not None and -1 not in timestamps_list:
+                        frame_timestamps = timestamps_list[idx_in_list]
+                    else:
+                        frame_timestamps = None
+                    sampled_fps = fps_list[idx_in_list] if fps_list is not None else None
+                    
+                    if frame_timestamps is not None:
+                        # Ensure lengths match (sometimes rounding might cause off-by-one if not careful, but usually safe here)
+                        if len(frame_timestamps) != len(num_of_tokens_list):
+                             logger.warning(f"Timestamp mismatch: {len(frame_timestamps)} vs {len(num_of_tokens_list)}")
+                             min_len = min(len(frame_timestamps), len(num_of_tokens_list))
+                             frame_timestamps = frame_timestamps[:min_len]
+                             num_of_tokens_list = num_of_tokens_list[:min_len]
+
+                        special_placeholder = [f"Frame-{i+1}-{frame_timestamps[i]:.2f}s: {self.image_start_token}{self.image_token * num_of_tokens}{self.image_end_token}" for i, num_of_tokens in enumerate(num_of_tokens_list)]
+                    else:
+                        special_placeholder = [f"Frame-{i+1}: {self.image_start_token}{self.image_token * num_of_tokens}{self.image_end_token}" for i, num_of_tokens in enumerate(num_of_tokens_list)]
+                    
+                    if sampled_fps is not None:
+                        special_placeholder = f"The {idx_mapper[idx_in_list]} video sampled with {sampled_fps:.2f} fps: " + "".join(special_placeholder)
+                    else:
+                        special_placeholder = f"The {idx_mapper[idx_in_list]} video: " + "".join(special_placeholder)
+                    
+                    unified_frame_list.append(video_inputs)
+                    num_of_videos_in_this_sample += 1
+                else:
+                    raise ValueError(f'Unknown media type: {media_type}')
+                return special_placeholder
+            return pattern.sub(repl, text)
+        
+        text = replace_in_text(text)
+        
+        if len(unified_frame_list) > 0:
+            # Concatenate all pixel values from all images/videos in this sample
+            pixel_values = torch.cat([frame['pixel_values'] for frame in unified_frame_list], dim=0)
+            # Concatenate grid hws
+            image_grid_hws = np.concatenate([frame['image_grid_hws'] for frame in unified_frame_list], axis=0)
+        else:
+            pixel_values = torch.empty(0)
+            image_grid_hws = np.empty(0)
+            
+        return text, pixel_values, image_grid_hws, num_of_images_in_this_sample, num_of_videos_in_this_sample
+    
+    def __call__(
+        self,
+        images: ImageInput = None,
+        text: Union[TextInput, PreTokenizedInput, List[TextInput], List[PreTokenizedInput]] = None,
+        audio=None,
+        videos: VideoInput = None,
+        **kwargs: Unpack[LocateAnythingProcessorKwargs],
+    ) -> BatchFeature:
+        output_kwargs = self._merge_kwargs(
+            LocateAnythingProcessorKwargs,
+            tokenizer_init_kwargs=self.tokenizer.init_kwargs,
+            **kwargs,
+        )
+        
+        if isinstance(text, str):
+            text_list = [text]
+        elif not isinstance(text, list) and not isinstance(text[0], str):
+            raise ValueError("Invalid input text. Please provide a string, or a list of strings")
+        elif isinstance(text, list) and isinstance(text[0], str):
+            text_list = text
+        
+        if images is None: images = []
+        if videos is None: videos = []
+        
+        pixel_values_list = []
+        image_grid_hws_list = []
+        new_sample_list = []
+        image_start_idx = 0
+        video_start_idx = 0
+        timestamps_batch = output_kwargs['videos_kwargs'].pop("timestamps", None)
+        fps_batch = output_kwargs['videos_kwargs'].pop("fps", None)
+        
+        for sample in text_list:
+            timestamps_list = timestamps_batch[video_start_idx:] if timestamps_batch is not None else None
+            fps_list = fps_batch[video_start_idx:] if fps_batch is not None else None
+            
+            sample, pixel_values, image_grid_hws, num_of_images_in_this_sample, num_of_videos_in_this_sample = self.replace_media_placeholder(
+                sample, images[image_start_idx:], videos[video_start_idx:], timestamps_list, fps_list, **output_kwargs
+            )
+            new_sample_list.append(sample)
+            
+            if pixel_values.numel() > 0:
+                pixel_values_list.append(pixel_values)
+                image_grid_hws_list.append(image_grid_hws)
+
+            image_start_idx += num_of_images_in_this_sample
+            video_start_idx += num_of_videos_in_this_sample
+
+        image_inputs = {}
+        if len(pixel_values_list) > 0:
+            # Concatenate across the batch
+            image_inputs['pixel_values'] = torch.cat(pixel_values_list, dim=0)
+            image_inputs['image_grid_hws'] = np.concatenate(image_grid_hws_list, axis=0)
+        
+        video_inputs = {} # Video data is merged into image_inputs now
+        text_inputs = self.tokenizer(new_sample_list, **output_kwargs["text_kwargs"])
+        
+        return BatchFeature(data={**text_inputs, **image_inputs, **video_inputs})
+
+    def batch_decode(self, *args, **kwargs):
+        return self.tokenizer.batch_decode(*args, **kwargs)
+
+    def decode(self, *args, **kwargs):
+        return self.tokenizer.decode(*args, **kwargs)
+
+    @property
+    def model_input_names(self):
+        tokenizer_input_names = self.tokenizer.model_input_names
+        image_processor_input_names = self.image_processor.model_input_names
+        return list(dict.fromkeys(tokenizer_input_names + image_processor_input_names))
+
+    def save_pretrained(self, save_directory, **kwargs):
+        if os.path.isfile(save_directory):
+            raise ValueError(f"Provided path ({save_directory}) should be a directory, not a file")
+        os.makedirs(save_directory, exist_ok=True)
+        outputs = super().save_pretrained(save_directory, **kwargs)
+        return outputs
+
+    @classmethod
+    def from_pretrained(cls, pretrained_model_name_or_path, **kwargs):
+        processor = super().from_pretrained(pretrained_model_name_or_path, **kwargs)
+        if isinstance(processor, tuple):
+            processor = processor[0]
+        return processor
+
+    def process_vision_info(
+        self,
+        conversations: list[dict] | list[list[dict]],
+        return_video_kwargs: bool = False,
+        video_reader_backend: str = "torchvision",
+    ) -> tuple[list[Image.Image] | None, list[torch.Tensor | list[Image.Image]] | None, Optional[dict]]:
+
+        vision_infos = self.extract_vision_info(conversations)
+        image_inputs = []
+        video_inputs = []
+        video_sample_fps_list = []
+        video_timestamps_list = []
+        
+        for vision_info in vision_infos:
+            if "image" in vision_info or "image_url" in vision_info:
+                image_inputs.append(fetch_image(vision_info))
+            elif "video" in vision_info:
+                video_input, video_sample_fps, video_timestamps = fetch_video(vision_info, return_video_sample_fps=True, video_reader_backend=video_reader_backend)
+                video_sample_fps_list.append(video_sample_fps)
+                video_inputs.append(video_input)
+                video_timestamps_list.append(video_timestamps)
+            else:
+                raise ValueError("image, image_url or video should in content.")
+        
+        if len(image_inputs) == 0:
+            image_inputs = None
+        if len(video_inputs) == 0:
+            video_inputs = None
+            
+        if return_video_kwargs:
+            return image_inputs, video_inputs, {'fps': video_sample_fps_list, 'timestamps': video_timestamps_list}
+        return image_inputs, video_inputs
+
+    def extract_vision_info(self, conversations: list[dict] | list[list[dict]]) -> list[dict]:
+        vision_infos = []
+        if isinstance(conversations[0], dict):
+            conversations = [conversations]
+        for conversation in conversations:
+            for message in conversation:
+                if isinstance(message["content"], list):
+                    for ele in message["content"]:
+                        if (
+                            "image" in ele
+                            or "image_url" in ele
+                            or "video" in ele
+                            or ele["type"] in ("image", "image_url", "video")
+                        ):
+                            vision_infos.append(ele)
+        return vision_infos
+    
+    def py_apply_chat_template(self, messages, tokenize=False, add_generation_prompt=False):
+        assert tokenize == False, "tokenize is not supported yet"
+        result = ""
+        image_count = 0
+        video_count = 0
+        
+        message_text = ""
+        for idx, message in enumerate(messages):
+            if message.get('role') != 'user': continue
+            content = message.get('content')
+            if isinstance(content, str):
+                message_text += content
+            elif isinstance(content, list):
+                for item in content:
+                    if isinstance(item, dict) and "text" in item:
+                        message_text += item["text"]
+                    elif isinstance(item, str):
+                        message_text += item
+                        
+        for idx, message in enumerate(messages):
+            if idx == 0 and message.get('role') != 'system':
+                result += "<|im_start|>system\n"
+                result += "You are a helpful assistant.\n"
+                result += "<|im_end|>\n"
+
+            result += f"<|im_start|>{message.get('role', '')}\n"
+            content = message.get('content')
+
+            if isinstance(content, str):
+                result += content
+                result += "<|im_end|>\n"
+            else:
+                for item in content:
+                    if (isinstance(item, dict) and (item.get('type') == 'image' or 'image' in item or 'image_url' in item)):
+                        image_count += 1
+                        candidate_token = f"<image-{image_count}>"
+                        if candidate_token not in message_text:
+                            result += candidate_token
+                    elif (isinstance(item, dict) and (item.get('type') == 'video' or 'video' in item)):
+                        video_count += 1
+                        candidate_token = f"<video-{video_count}>"
+                        if candidate_token not in message_text:
+                            result += candidate_token
+                    elif isinstance(item, dict) and 'text' in item:
+                        result += item['text']
+                    elif isinstance(item, str):
+                        result += item
+                result += "<|im_end|>\n"
+
+        if add_generation_prompt:
+            result += "<|im_start|>assistant\n"
+
+        return result
+
+
+    @classmethod
+    def from_args_and_dict(cls, args, processor_dict: dict[str, Any], **kwargs):
+        processor_dict = processor_dict.copy()
+        return_unused_kwargs = kwargs.pop("return_unused_kwargs", False)
+
+        if "processor_class" in processor_dict:
+            del processor_dict["processor_class"]
+
+        unused_kwargs = cls.validate_init_kwargs(processor_config=processor_dict, valid_kwargs=cls.valid_kwargs)
+        processor = cls(*args, **processor_dict)
+
+        for key in set(kwargs.keys()):
+            if hasattr(processor, key):
+                setattr(processor, key, kwargs.pop(key))
+
+        if isinstance(unused_kwargs, dict):
+            kwargs.update(unused_kwargs)
+        logger.info(f"Processor {processor}")
+        if return_unused_kwargs:
+            return processor, kwargs
+        else:
+            return processor
+        
+        
+__all__ = ["LocateAnythingProcessor"]

processor_config.json

@@ -0,0 +1,18 @@
+{
+  "auto_map": {
+    "AutoImageProcessor": "image_processing_locateanything.LocateAnythingImageProcessor",
+    "AutoProcessor": "processing_locateanything.LocateAnythingProcessor"
+  },
+  "image_end_token": "</img>",
+  "image_placeholder": "image",
+  "image_start_token": "<img>",
+  "image_token": "<IMG_CONTEXT>",
+  "merge_kernel_size": [
+    2,
+    2
+  ],
+  "processor_class": "LocateAnythingProcessor",
+  "video_placeholder": "video",
+  "video_token": "<IMG_CONTEXT>",
+  "patch_size": 14
+}

pyproject.toml

@@ -0,0 +1,18 @@
+[build-system]
+requires = ["setuptools>=68", "wheel"]
+build-backend = "setuptools.build_meta"
+
+[project]
+name = "locateanything-la-flash"
+version = "0.1.0"
+description = "LocateAnything batch utils with LA Flash inference backend."
+requires-python = ">=3.10"
+dependencies = [
+  "numpy",
+  "pillow",
+  "torch",
+  "transformers",
+]
+
+[tool.setuptools]
+packages = ["batch_utils", "kernel_utils"]

special_tokens_map.json

@@ -0,0 +1,1053 @@
+{
+  "additional_special_tokens": [
+    "<|im_start|>",
+    "<|im_end|>",
+    "<|object_ref_start|>",
+    "<|object_ref_end|>",
+    "<|box_start|>",
+    "<|box_end|>",
+    "<|quad_start|>",
+    "<|quad_end|>",
+    "<|vision_start|>",
+    "<|vision_end|>",
+    "<|vision_pad|>",
+    "<|image_pad|>",
+    "<|video_pad|>",
+    "<IMG_CONTEXT>",
+    "<img>",
+    "</img>",
+    "<box>",
+    "</box>",
+    "<quad>",
+    "</quad>",
+    "<ref>",
+    "</ref>",
+    "<interval>",
+    "</interval>",
+    "<text_mask>",
+    "<0>",
+    "<1>",
+    "<2>",
+    "<3>",
+    "<4>",
+    "<5>",
+    "<6>",
+    "<7>",
+    "<8>",
+    "<9>",
+    "<10>",
+    "<11>",
+    "<12>",
+    "<13>",
+    "<14>",
+    "<15>",
+    "<16>",
+    "<17>",
+    "<18>",
+    "<19>",
+    "<20>",
+    "<21>",
+    "<22>",
+    "<23>",
+    "<24>",
+    "<25>",
+    "<26>",
+    "<27>",
+    "<28>",
+    "<29>",
+    "<30>",
+    "<31>",
+    "<32>",
+    "<33>",
+    "<34>",
+    "<35>",
+    "<36>",
+    "<37>",
+    "<38>",
+    "<39>",
+    "<40>",
+    "<41>",
+    "<42>",
+    "<43>",
+    "<44>",
+    "<45>",
+    "<46>",
+    "<47>",
+    "<48>",
+    "<49>",
+    "<50>",
+    "<51>",
+    "<52>",
+    "<53>",
+    "<54>",
+    "<55>",
+    "<56>",
+    "<57>",
+    "<58>",
+    "<59>",
+    "<60>",
+    "<61>",
+    "<62>",
+    "<63>",
+    "<64>",
+    "<65>",
+    "<66>",
+    "<67>",
+    "<68>",
+    "<69>",
+    "<70>",
+    "<71>",
+    "<72>",
+    "<73>",
+    "<74>",
+    "<75>",
+    "<76>",
+    "<77>",
+    "<78>",
+    "<79>",
+    "<80>",
+    "<81>",
+    "<82>",
+    "<83>",
+    "<84>",
+    "<85>",
+    "<86>",
+    "<87>",
+    "<88>",
+    "<89>",
+    "<90>",
+    "<91>",
+    "<92>",
+    "<93>",
+    "<94>",
+    "<95>",
+    "<96>",
+    "<97>",
+    "<98>",
+    "<99>",
+    "<100>",
+    "<101>",
+    "<102>",
+    "<103>",
+    "<104>",
+    "<105>",
+    "<106>",
+    "<107>",
+    "<108>",
+    "<109>",
+    "<110>",
+    "<111>",
+    "<112>",
+    "<113>",
+    "<114>",
+    "<115>",
+    "<116>",
+    "<117>",
+    "<118>",
+    "<119>",
+    "<120>",
+    "<121>",
+    "<122>",
+    "<123>",
+    "<124>",
+    "<125>",
+    "<126>",
+    "<127>",
+    "<128>",
+    "<129>",
+    "<130>",
+    "<131>",
+    "<132>",
+    "<133>",
+    "<134>",
+    "<135>",
+    "<136>",
+    "<137>",
+    "<138>",
+    "<139>",
+    "<140>",
+    "<141>",
+    "<142>",
+    "<143>",
+    "<144>",
+    "<145>",
+    "<146>",
+    "<147>",
+    "<148>",
+    "<149>",
+    "<150>",
+    "<151>",
+    "<152>",
+    "<153>",
+    "<154>",
+    "<155>",
+    "<156>",
+    "<157>",
+    "<158>",
+    "<159>",
+    "<160>",
+    "<161>",
+    "<162>",
+    "<163>",
+    "<164>",
+    "<165>",
+    "<166>",
+    "<167>",
+    "<168>",
+    "<169>",
+    "<170>",
+    "<171>",
+    "<172>",
+    "<173>",
+    "<174>",
+    "<175>",
+    "<176>",
+    "<177>",
+    "<178>",
+    "<179>",
+    "<180>",
+    "<181>",
+    "<182>",
+    "<183>",
+    "<184>",
+    "<185>",
+    "<186>",
+    "<187>",
+    "<188>",
+    "<189>",
+    "<190>",
+    "<191>",
+    "<192>",
+    "<193>",
+    "<194>",
+    "<195>",
+    "<196>",
+    "<197>",
+    "<198>",
+    "<199>",
+    "<200>",
+    "<201>",
+    "<202>",
+    "<203>",
+    "<204>",
+    "<205>",
+    "<206>",
+    "<207>",
+    "<208>",
+    "<209>",
+    "<210>",
+    "<211>",
+    "<212>",
+    "<213>",
+    "<214>",
+    "<215>",
+    "<216>",
+    "<217>",
+    "<218>",
+    "<219>",
+    "<220>",
+    "<221>",
+    "<222>",
+    "<223>",
+    "<224>",
+    "<225>",
+    "<226>",
+    "<227>",
+    "<228>",
+    "<229>",
+    "<230>",
+    "<231>",
+    "<232>",
+    "<233>",
+    "<234>",
+    "<235>",
+    "<236>",
+    "<237>",
+    "<238>",
+    "<239>",
+    "<240>",
+    "<241>",
+    "<242>",
+    "<243>",
+    "<244>",
+    "<245>",
+    "<246>",
+    "<247>",
+    "<248>",
+    "<249>",
+    "<250>",
+    "<251>",
+    "<252>",
+    "<253>",
+    "<254>",
+    "<255>",
+    "<256>",
+    "<257>",
+    "<258>",
+    "<259>",
+    "<260>",
+    "<261>",
+    "<262>",
+    "<263>",
+    "<264>",
+    "<265>",
+    "<266>",
+    "<267>",
+    "<268>",
+    "<269>",
+    "<270>",
+    "<271>",
+    "<272>",
+    "<273>",
+    "<274>",
+    "<275>",
+    "<276>",
+    "<277>",
+    "<278>",
+    "<279>",
+    "<280>",
+    "<281>",
+    "<282>",
+    "<283>",
+    "<284>",
+    "<285>",
+    "<286>",
+    "<287>",
+    "<288>",
+    "<289>",
+    "<290>",
+    "<291>",
+    "<292>",
+    "<293>",
+    "<294>",
+    "<295>",
+    "<296>",
+    "<297>",
+    "<298>",
+    "<299>",
+    "<300>",
+    "<301>",
+    "<302>",
+    "<303>",
+    "<304>",
+    "<305>",
+    "<306>",
+    "<307>",
+    "<308>",
+    "<309>",
+    "<310>",
+    "<311>",
+    "<312>",
+    "<313>",
+    "<314>",
+    "<315>",
+    "<316>",
+    "<317>",
+    "<318>",
+    "<319>",
+    "<320>",
+    "<321>",
+    "<322>",
+    "<323>",
+    "<324>",
+    "<325>",
+    "<326>",
+    "<327>",
+    "<328>",
+    "<329>",
+    "<330>",
+    "<331>",
+    "<332>",
+    "<333>",
+    "<334>",
+    "<335>",
+    "<336>",
+    "<337>",
+    "<338>",
+    "<339>",
+    "<340>",
+    "<341>",
+    "<342>",
+    "<343>",
+    "<344>",
+    "<345>",
+    "<346>",
+    "<347>",
+    "<348>",
+    "<349>",
+    "<350>",
+    "<351>",
+    "<352>",
+    "<353>",
+    "<354>",
+    "<355>",
+    "<356>",
+    "<357>",
+    "<358>",
+    "<359>",
+    "<360>",
+    "<361>",
+    "<362>",
+    "<363>",
+    "<364>",
+    "<365>",
+    "<366>",
+    "<367>",
+    "<368>",
+    "<369>",
+    "<370>",
+    "<371>",
+    "<372>",
+    "<373>",
+    "<374>",
+    "<375>",
+    "<376>",
+    "<377>",
+    "<378>",
+    "<379>",
+    "<380>",
+    "<381>",
+    "<382>",
+    "<383>",
+    "<384>",
+    "<385>",
+    "<386>",
+    "<387>",
+    "<388>",
+    "<389>",
+    "<390>",
+    "<391>",
+    "<392>",
+    "<393>",
+    "<394>",
+    "<395>",
+    "<396>",
+    "<397>",
+    "<398>",
+    "<399>",
+    "<400>",
+    "<401>",
+    "<402>",
+    "<403>",
+    "<404>",
+    "<405>",
+    "<406>",
+    "<407>",
+    "<408>",
+    "<409>",
+    "<410>",
+    "<411>",
+    "<412>",
+    "<413>",
+    "<414>",
+    "<415>",
+    "<416>",
+    "<417>",
+    "<418>",
+    "<419>",
+    "<420>",
+    "<421>",
+    "<422>",
+    "<423>",
+    "<424>",
+    "<425>",
+    "<426>",
+    "<427>",
+    "<428>",
+    "<429>",
+    "<430>",
+    "<431>",
+    "<432>",
+    "<433>",
+    "<434>",
+    "<435>",
+    "<436>",
+    "<437>",
+    "<438>",
+    "<439>",
+    "<440>",
+    "<441>",
+    "<442>",
+    "<443>",
+    "<444>",
+    "<445>",
+    "<446>",
+    "<447>",
+    "<448>",
+    "<449>",
+    "<450>",
+    "<451>",
+    "<452>",
+    "<453>",
+    "<454>",
+    "<455>",
+    "<456>",
+    "<457>",
+    "<458>",
+    "<459>",
+    "<460>",
+    "<461>",
+    "<462>",
+    "<463>",
+    "<464>",
+    "<465>",
+    "<466>",
+    "<467>",
+    "<468>",
+    "<469>",
+    "<470>",
+    "<471>",
+    "<472>",
+    "<473>",
+    "<474>",
+    "<475>",
+    "<476>",
+    "<477>",
+    "<478>",
+    "<479>",
+    "<480>",
+    "<481>",
+    "<482>",
+    "<483>",
+    "<484>",
+    "<485>",
+    "<486>",
+    "<487>",
+    "<488>",
+    "<489>",
+    "<490>",
+    "<491>",
+    "<492>",
+    "<493>",
+    "<494>",
+    "<495>",
+    "<496>",
+    "<497>",
+    "<498>",
+    "<499>",
+    "<500>",
+    "<501>",
+    "<502>",
+    "<503>",
+    "<504>",
+    "<505>",
+    "<506>",
+    "<507>",
+    "<508>",
+    "<509>",
+    "<510>",
+    "<511>",
+    "<512>",
+    "<513>",
+    "<514>",
+    "<515>",
+    "<516>",
+    "<517>",
+    "<518>",
+    "<519>",
+    "<520>",
+    "<521>",
+    "<522>",
+    "<523>",
+    "<524>",
+    "<525>",
+    "<526>",
+    "<527>",
+    "<528>",
+    "<529>",
+    "<530>",
+    "<531>",
+    "<532>",
+    "<533>",
+    "<534>",
+    "<535>",
+    "<536>",
+    "<537>",
+    "<538>",
+    "<539>",
+    "<540>",
+    "<541>",
+    "<542>",
+    "<543>",
+    "<544>",
+    "<545>",
+    "<546>",
+    "<547>",
+    "<548>",
+    "<549>",
+    "<550>",
+    "<551>",
+    "<552>",
+    "<553>",
+    "<554>",
+    "<555>",
+    "<556>",
+    "<557>",
+    "<558>",
+    "<559>",
+    "<560>",
+    "<561>",
+    "<562>",
+    "<563>",
+    "<564>",
+    "<565>",
+    "<566>",
+    "<567>",
+    "<568>",
+    "<569>",
+    "<570>",
+    "<571>",
+    "<572>",
+    "<573>",
+    "<574>",
+    "<575>",
+    "<576>",
+    "<577>",
+    "<578>",
+    "<579>",
+    "<580>",
+    "<581>",
+    "<582>",
+    "<583>",
+    "<584>",
+    "<585>",
+    "<586>",
+    "<587>",
+    "<588>",
+    "<589>",
+    "<590>",
+    "<591>",
+    "<592>",
+    "<593>",
+    "<594>",
+    "<595>",
+    "<596>",
+    "<597>",
+    "<598>",
+    "<599>",
+    "<600>",
+    "<601>",
+    "<602>",
+    "<603>",
+    "<604>",
+    "<605>",
+    "<606>",
+    "<607>",
+    "<608>",
+    "<609>",
+    "<610>",
+    "<611>",
+    "<612>",
+    "<613>",
+    "<614>",
+    "<615>",
+    "<616>",
+    "<617>",
+    "<618>",
+    "<619>",
+    "<620>",
+    "<621>",
+    "<622>",
+    "<623>",
+    "<624>",
+    "<625>",
+    "<626>",
+    "<627>",
+    "<628>",
+    "<629>",
+    "<630>",
+    "<631>",
+    "<632>",
+    "<633>",
+    "<634>",
+    "<635>",
+    "<636>",
+    "<637>",
+    "<638>",
+    "<639>",
+    "<640>",
+    "<641>",
+    "<642>",
+    "<643>",
+    "<644>",
+    "<645>",
+    "<646>",
+    "<647>",
+    "<648>",
+    "<649>",
+    "<650>",
+    "<651>",
+    "<652>",
+    "<653>",
+    "<654>",
+    "<655>",
+    "<656>",
+    "<657>",
+    "<658>",
+    "<659>",
+    "<660>",
+    "<661>",
+    "<662>",
+    "<663>",
+    "<664>",
+    "<665>",
+    "<666>",
+    "<667>",
+    "<668>",
+    "<669>",
+    "<670>",
+    "<671>",
+    "<672>",
+    "<673>",
+    "<674>",
+    "<675>",
+    "<676>",
+    "<677>",
+    "<678>",
+    "<679>",
+    "<680>",
+    "<681>",
+    "<682>",
+    "<683>",
+    "<684>",
+    "<685>",
+    "<686>",
+    "<687>",
+    "<688>",
+    "<689>",
+    "<690>",
+    "<691>",
+    "<692>",
+    "<693>",
+    "<694>",
+    "<695>",
+    "<696>",
+    "<697>",
+    "<698>",
+    "<699>",
+    "<700>",
+    "<701>",
+    "<702>",
+    "<703>",
+    "<704>",
+    "<705>",
+    "<706>",
+    "<707>",
+    "<708>",
+    "<709>",
+    "<710>",
+    "<711>",
+    "<712>",
+    "<713>",
+    "<714>",
+    "<715>",
+    "<716>",
+    "<717>",
+    "<718>",
+    "<719>",
+    "<720>",
+    "<721>",
+    "<722>",
+    "<723>",
+    "<724>",
+    "<725>",
+    "<726>",
+    "<727>",
+    "<728>",
+    "<729>",
+    "<730>",
+    "<731>",
+    "<732>",
+    "<733>",
+    "<734>",
+    "<735>",
+    "<736>",
+    "<737>",
+    "<738>",
+    "<739>",
+    "<740>",
+    "<741>",
+    "<742>",
+    "<743>",
+    "<744>",
+    "<745>",
+    "<746>",
+    "<747>",
+    "<748>",
+    "<749>",
+    "<750>",
+    "<751>",
+    "<752>",
+    "<753>",
+    "<754>",
+    "<755>",
+    "<756>",
+    "<757>",
+    "<758>",
+    "<759>",
+    "<760>",
+    "<761>",
+    "<762>",
+    "<763>",
+    "<764>",
+    "<765>",
+    "<766>",
+    "<767>",
+    "<768>",
+    "<769>",
+    "<770>",
+    "<771>",
+    "<772>",
+    "<773>",
+    "<774>",
+    "<775>",
+    "<776>",
+    "<777>",
+    "<778>",
+    "<779>",
+    "<780>",
+    "<781>",
+    "<782>",
+    "<783>",
+    "<784>",
+    "<785>",
+    "<786>",
+    "<787>",
+    "<788>",
+    "<789>",
+    "<790>",
+    "<791>",
+    "<792>",
+    "<793>",
+    "<794>",
+    "<795>",
+    "<796>",
+    "<797>",
+    "<798>",
+    "<799>",
+    "<800>",
+    "<801>",
+    "<802>",
+    "<803>",
+    "<804>",
+    "<805>",
+    "<806>",
+    "<807>",
+    "<808>",
+    "<809>",
+    "<810>",
+    "<811>",
+    "<812>",
+    "<813>",
+    "<814>",
+    "<815>",
+    "<816>",
+    "<817>",
+    "<818>",
+    "<819>",
+    "<820>",
+    "<821>",
+    "<822>",
+    "<823>",
+    "<824>",
+    "<825>",
+    "<826>",
+    "<827>",
+    "<828>",
+    "<829>",
+    "<830>",
+    "<831>",
+    "<832>",
+    "<833>",
+    "<834>",
+    "<835>",
+    "<836>",
+    "<837>",
+    "<838>",
+    "<839>",
+    "<840>",
+    "<841>",
+    "<842>",
+    "<843>",
+    "<844>",
+    "<845>",
+    "<846>",
+    "<847>",
+    "<848>",
+    "<849>",
+    "<850>",
+    "<851>",
+    "<852>",
+    "<853>",
+    "<854>",
+    "<855>",
+    "<856>",
+    "<857>",
+    "<858>",
+    "<859>",
+    "<860>",
+    "<861>",
+    "<862>",
+    "<863>",
+    "<864>",
+    "<865>",
+    "<866>",
+    "<867>",
+    "<868>",
+    "<869>",
+    "<870>",
+    "<871>",
+    "<872>",
+    "<873>",
+    "<874>",
+    "<875>",
+    "<876>",
+    "<877>",
+    "<878>",
+    "<879>",
+    "<880>",
+    "<881>",
+    "<882>",
+    "<883>",
+    "<884>",
+    "<885>",
+    "<886>",
+    "<887>",
+    "<888>",
+    "<889>",
+    "<890>",
+    "<891>",
+    "<892>",
+    "<893>",
+    "<894>",
+    "<895>",
+    "<896>",
+    "<897>",
+    "<898>",
+    "<899>",
+    "<900>",
+    "<901>",
+    "<902>",
+    "<903>",
+    "<904>",
+    "<905>",
+    "<906>",
+    "<907>",
+    "<908>",
+    "<909>",
+    "<910>",
+    "<911>",
+    "<912>",
+    "<913>",
+    "<914>",
+    "<915>",
+    "<916>",
+    "<917>",
+    "<918>",
+    "<919>",
+    "<920>",
+    "<921>",
+    "<922>",
+    "<923>",
+    "<924>",
+    "<925>",
+    "<926>",
+    "<927>",
+    "<928>",
+    "<929>",
+    "<930>",
+    "<931>",
+    "<932>",
+    "<933>",
+    "<934>",
+    "<935>",
+    "<936>",
+    "<937>",
+    "<938>",
+    "<939>",
+    "<940>",
+    "<941>",
+    "<942>",
+    "<943>",
+    "<944>",
+    "<945>",
+    "<946>",
+    "<947>",
+    "<948>",
+    "<949>",
+    "<950>",
+    "<951>",
+    "<952>",
+    "<953>",
+    "<954>",
+    "<955>",
+    "<956>",
+    "<957>",
+    "<958>",
+    "<959>",
+    "<960>",
+    "<961>",
+    "<962>",
+    "<963>",
+    "<964>",
+    "<965>",
+    "<966>",
+    "<967>",
+    "<968>",
+    "<969>",
+    "<970>",
+    "<971>",
+    "<972>",
+    "<973>",
+    "<974>",
+    "<975>",
+    "<976>",
+    "<977>",
+    "<978>",
+    "<979>",
+    "<980>",
+    "<981>",
+    "<982>",
+    "<983>",
+    "<984>",
+    "<985>",
+    "<986>",
+    "<987>",
+    "<988>",
+    "<989>",
+    "<990>",
+    "<991>",
+    "<992>",
+    "<993>",
+    "<994>",
+    "<995>",
+    "<996>",
+    "<997>",
+    "<998>",
+    "<999>",
+    "<1000>",
+    "<null>",
+    "<switch>",
+    {
+      "content": "</c>",
+      "lstrip": false,
+      "normalized": false,
+      "rstrip": false,
+      "single_word": false
+    }
+  ],
+  "eos_token": {
+    "content": "<|im_end|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  },
+  "pad_token": {
+    "content": "<|endoftext|>",
+    "lstrip": false,
+    "normalized": false,
+    "rstrip": false,
+    "single_word": false
+  }
+}

training_args.bin

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:995b5f0a2fe72453ddc8ce97e1a93747554ec3ec0ac92d86e82a57050db51b85
+size 7288