update model

LICENSE

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

generative_r2d2_fast.py

@@ -10,9 +10,10 @@ from dataclasses import dataclass
 from typing import Optional
 from transformers import PreTrainedModel, PretrainedConfig, GPT2Config
 from transformers.modeling_outputs import MaskedLMOutput
-from lmup.gpst.gpt2_flash_attn import GPT2Model
-#from lmup.gpst.config import *
-from lmup.gpst.r2d2_insideoutside import *
+from gpst.gpt2_flash_attn import GPT2Model
+from gpst.backend_loader import load_cpp_backend
+cppbackend = load_cpp_backend()
+from gpst.r2d2_insideoutside import *
 import copy
 
 def load_model(model, model_path, strict=True):
@@ -186,46 +187,9 @@ class FastGenerativeR2D2(nn.Module):
         next_token_indices = next_token_indices[:, :max_input_len + 1]
         return next_token_indices, chunk_input_ids
 
-    def gpst_prep(self, input_ids, attention_mask, max_len=96):
-        """
-        - `input_ids`: token ids, with padding id 0, such that 2048 tokens are reached (example shape: 66x81) for 66 sentences with max length 81
-        - `chunk_input_ids`: The same tokens, but without padding
-        - `chunk_masks`: Shape 2x1024, and in each row values 1-something to indicate the sentence in `input_ids`. Starting with 1!
-        - `masks`: Same shape as the input ids, 1 for input tokens, 0 for padding
-        - `group_ids`: One-dim with length 66 (num sents), values 0 and 1 to indicate which sentence goes in which row of the `chunk_input_ids` and `chunk_masks`
-        - `atom_spans` is None, `span_ids` is an empty list, `external_vocab_ids` is None
-        """
-        device = attention_mask.device
-        max_batch_len = attention_mask.sum(dim=1).max().item() # old for formal language, depends on pad_token_id==0 batch["input_ids"].nonzero()[:,1].max()
-        #chunk_input_ids = batch["input_ids"].view(-1) # old version, removes padding
-        #chunk_input_ids = chunk_input_ids[chunk_input_ids != 0].unsqueeze(0)
-        input_ids = input_ids[:,:max_batch_len]
-        attention_mask = attention_mask[:,:max_batch_len]
-        chunk_input_ids = input_ids # when the assertion is true
-        #group_ids = model_inputs["input_ids"].nonzero()[:,0]
-        group_ids = torch.arange(input_ids.shape[0], dtype=int, device=device)
-        #chunk_masks = (batch["input_ids"].nonzero()[:,0].unsqueeze(0))+1
-        chunk_masks = attention_mask * (group_ids+1).unsqueeze(1)
-        gpst_batch = {
-            "input_ids": input_ids,        # batch_size (sents) x max_seq_length
-            "masks": attention_mask,       
-            "chunk_input_ids": chunk_input_ids,     # 
-            "group_ids": group_ids.cpu().numpy(),
-            "chunk_masks": chunk_masks
-        }
-        return gpst_batch
-
     def forward(self, chunk_input_ids= None, chunk_masks=None, input_ids=None, masks=None, eos_labels=None, group_ids=None, 
-                atom_spans=None, span_ids=None, external_vocab_ids=None, attention_mask=None,
+                atom_spans=None, span_ids=None, external_vocab_ids=None, 
                 coeff=1.0, temperature=1.0, past_key_values=None):
-        
-        gpst_batch = self.gpst_prep(input_ids, attention_mask)
-        input_ids = gpst_batch["input_ids"]
-        masks = gpst_batch["masks"]
-        chunk_input_ids = gpst_batch["chunk_input_ids"]
-        group_ids = gpst_batch["group_ids"]
-        chunk_masks = gpst_batch["chunk_masks"]
-        
         batch_size = max(group_ids) + 1
         r2d2_input_ids = torch.where(chunk_input_ids == -100, 0, chunk_input_ids)
         input_embeddings = self.embeddings(r2d2_input_ids)

gpst/Llama_flash_attn.py

@@ -0,0 +1,284 @@
+# Copyright (c) Meta Platforms, Inc. and affiliates.
+# This software may be used and distributed according to the terms of the GNU General Public License version 3.
+
+from typing import Optional, Tuple, Type
+from dataclasses import dataclass
+import math
+
+import torch
+from torch import nn
+import torch.nn.functional as F
+
+
+@dataclass
+class ModelArgs:
+    dim: int = 512
+    n_layers: int = 8
+    n_heads: int = 8
+    vocab_size: int = -1  # defined later by tokenizer
+    multiple_of: int = 256  # make SwiGLU hidden layer size multiple of large power of 2
+    norm_eps: float = 1e-5
+
+    max_seq_len: int = 1024
+    apply_norm: bool = True
+
+
+class RMSNorm(torch.nn.Module):
+    def __init__(self, dim: int, eps: float = 1e-6):
+        super().__init__()
+        self.eps = eps
+        self.weight = nn.Parameter(torch.ones(dim))
+
+    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.eps)
+        return self.weight * hidden_states.to(input_dtype)
+
+
+# def precompute_freqs_cis(dim: int, end: int, theta: float = 10000.0):
+#     freqs = 1.0 / (theta ** (torch.arange(0, dim, 2)[: (dim // 2)] / dim))
+#     t = torch.arange(end, device=freqs.device)  # type: ignore
+#     freqs = torch.outer(t, freqs)  # type: ignore
+#     freqs_cis = torch.polar(torch.ones_like(freqs), freqs)  # complex64
+#     return freqs_cis
+
+
+def reshape_for_broadcast(freqs_cis: torch.Tensor, x: torch.Tensor):
+    ndim = x.ndim
+    assert 0 <= 1 < ndim
+    assert freqs_cis.shape == (x.shape[1], x.shape[-1]), f'{freqs_cis.shape},{(x.shape[1], x.shape[-1])}'
+    shape = [d if i == 1 or i == ndim - 1 else 1 for i, d in enumerate(x.shape)]
+    return freqs_cis.view(*shape)
+
+
+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)
+
+def apply_rotary_emb(
+    xq: torch.Tensor,
+    xk: torch.Tensor,
+    freqs_cis: torch.Tensor,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    # xq_ = torch.view_as_complex(xq.reshape(*xq.shape[:-1], -1, 2))
+    # xk_ = torch.view_as_complex(xk.reshape(*xk.shape[:-1], -1, 2))
+    # freqs_cis = reshape_for_broadcast(freqs_cis, xq_)
+    # xq_out = torch.view_as_real(xq_ * freqs_cis).flatten(3)
+    # xk_out = torch.view_as_real(xk_ * freqs_cis).flatten(3)
+    # return xq_out.type_as(xq), xk_out.type_as(xk)
+
+    cos, sin = freqs_cis
+    cos = cos.unsqueeze(1)
+    sin = sin.unsqueeze(1)
+    q_embed = (xq * cos) + (rotate_half(xq) * sin)
+    k_embed = (xk * cos) + (rotate_half(xk) * sin)
+    return q_embed, k_embed
+
+
+class Attention(nn.Module):
+    def __init__(self, args: ModelArgs):
+        super().__init__()
+
+        self.n_local_heads = args.n_heads
+        self.head_dim = args.dim // args.n_heads
+
+        self.wq = nn.Linear(
+            args.dim,
+            args.n_heads * self.head_dim,
+            bias=False,
+        )
+        self.wk = nn.Linear(
+            args.dim,
+            args.n_heads * self.head_dim,
+            bias=False,
+        )
+        self.wv = nn.Linear(
+            args.dim,
+            args.n_heads * self.head_dim,
+            bias=False,
+        )
+        self.wo = nn.Linear(
+            args.dim,
+            args.n_heads * self.head_dim,
+            bias=False,
+        )
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        freqs_cis: torch.Tensor,
+        attn_mask: Optional[torch.Tensor],
+    ):
+        bsz, seqlen, _ = x.shape
+        xq, xk, xv = self.wq(x), self.wk(x), self.wv(x)
+
+        xq = xq.view(bsz, seqlen, self.n_local_heads, self.head_dim)
+        xk = xk.view(bsz, seqlen, self.n_local_heads, self.head_dim)
+        xv = xv.view(bsz, seqlen, self.n_local_heads, self.head_dim)
+
+        xq = xq.transpose(1, 2)
+        xk = xk.transpose(1, 2)
+        values = xv.transpose(1, 2)
+
+        xq, xk = apply_rotary_emb(xq, xk, freqs_cis=freqs_cis)
+        if True:
+            #with torch.backends.cuda.sdpa_kernel():
+            if attn_mask is not None:
+                if len(attn_mask.shape) == 2:
+                    output = F.scaled_dot_product_attention(xq, xk, values, attn_mask=attn_mask.unsqueeze(1).unsqueeze(1))
+                elif len(attn_mask.shape) == 3:
+                    output = F.scaled_dot_product_attention(xq, xk, values, attn_mask=attn_mask.unsqueeze(1).to(xq.dtype))
+                else:
+                    raise Exception("Unsupported mask format")
+            else:
+                output = F.scaled_dot_product_attention(xq, xk, values, is_causal=True)
+
+        # L, S = xq.size(-2), keys.size(-2)
+        # scale_factor = 1 / math.sqrt(xq.size(-1))
+        # attn_weight = xq @ keys.transpose(-2, -1) * scale_factor
+        # attn_weight += attn_mask.unsqueeze(1).unsqueeze(1)
+        # attn_weight = torch.softmax(attn_weight, dim=-1)
+        # output = attn_weight @ values
+
+        output = output.transpose(1, 2).reshape(bsz, seqlen, -1)
+
+        return self.wo(output)
+
+
+class FeedForward(nn.Module):
+    def __init__(
+        self,
+        dim: int,
+        hidden_dim: int,
+        multiple_of: int,
+    ):
+        super().__init__()
+        hidden_dim = int(2 * hidden_dim / 3)
+        hidden_dim = multiple_of * ((hidden_dim + multiple_of - 1) // multiple_of)
+
+        self.w1 = nn.Linear(dim, hidden_dim, bias=False)
+        self.w2 = nn.Linear(
+            hidden_dim,
+            dim,
+            bias=False,
+        )
+        self.w3 = nn.Linear(
+            dim,
+            hidden_dim,
+            bias=False,
+        )
+
+    def forward(self, x):
+        return self.w2(F.silu(self.w1(x)) * self.w3(x))
+
+
+class TransformerBlock(nn.Module):
+    def __init__(self, layer_id: int, args: ModelArgs):
+        super().__init__()
+        self.n_heads = args.n_heads
+        self.dim = args.dim
+        self.head_dim = args.dim // args.n_heads
+        self.attention = Attention(args)
+        self.feed_forward = FeedForward(
+            dim=args.dim, hidden_dim=4 * args.dim, multiple_of=args.multiple_of
+        )
+        self.layer_id = layer_id
+        self.attention_norm = RMSNorm(args.dim, eps=args.norm_eps)
+        self.ffn_norm = RMSNorm(args.dim, eps=args.norm_eps)
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        freqs_cis: torch.Tensor,
+        mask: Optional[torch.Tensor],
+    ):
+        h = x + self.attention.forward(
+            self.attention_norm(x), freqs_cis, mask
+        )
+        out = h + self.feed_forward.forward(self.ffn_norm(h))
+        return out
+
+
+def convert_linear_to_bnb(float_linear):
+    new_layer = InferenceQuantizedLinear(
+        float_linear.in_features,
+        float_linear.out_features,
+        bias=float_linear.bias is not None,
+    )
+    new_layer._parameters["weight"] = bnb.nn.Int8Params(
+        float_linear.weight.data.cpu(),
+        requires_grad=False,
+        has_fp16_weights=False,
+    )
+    if float_linear.bias is not None:
+        new_layer._parameters["bias"] = float_linear.bias
+    return new_layer
+
+
+class Transformer(nn.Module):
+    def __init__(self, params: ModelArgs):
+        super().__init__()
+        self.params = params
+        self.vocab_size = params.vocab_size
+        self.n_layers = params.n_layers
+
+        self.tok_embeddings = torch.nn.Embedding(params.vocab_size, params.dim)
+
+        self.layers = torch.nn.ModuleList()
+        for layer_id in range(params.n_layers):
+            self.layers.append(TransformerBlock(layer_id, params))
+
+        if params.apply_norm:
+            self.norm = RMSNorm(params.dim, eps=params.norm_eps)
+        else:
+            self.norm = lambda x: x
+
+        # self._freqs_cis = None
+        self._init_sin_cos()
+    
+    def _init_sin_cos(self, base=10000.0):
+        device=next(self.parameters()).device
+        dim = self.params.dim // self.params.n_heads
+        inv_freq = 1.0 / (base ** (torch.arange(0, dim, 2).float().to(device) / dim))
+        self.register_buffer("inv_freq", inv_freq, persistent=False)
+
+        t = torch.arange(self.params.max_seq_len, device=device, dtype=self.inv_freq.dtype)
+
+        freqs = torch.einsum("i,j->ij", 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)
+        dtype=torch.get_default_dtype()
+        self.register_buffer("cos_cached", emb.cos().to(dtype), persistent=False)
+        self.register_buffer("sin_cached", emb.sin().to(dtype), persistent=False)
+        
+    def freqs_cis(self, position_ids, h):
+        # if self._freqs_cis is None:
+            # self._freqs_cis = precompute_freqs_cis(
+            #     self.params.dim // self.params.n_heads, self.params.max_seq_len
+            # )
+            # self._freqs_cis = self._freqs_cis.to(next(self.parameters()).device)
+           
+        # return self._freqs_cis
+
+        return (
+            self.cos_cached[position_ids].to(dtype=h.dtype),
+            self.sin_cached[position_ids].to(dtype=h.dtype),
+        )
+
+    def forward(self, tokens: torch.Tensor, attn_mask: torch.Tensor, position_ids: torch.Tensor):
+        _bsz, seqlen = tokens.shape
+        h = self.tok_embeddings(tokens)
+        if len(position_ids.shape) == 1:
+            position_ids = position_ids.unsqueeze(0).repeat(_bsz, 1)
+        freqs_cis = self.freqs_cis(position_ids, h)
+
+        for layer in self.layers:
+            h = layer(h, freqs_cis, mask=attn_mask)
+
+        h = self.norm(h)
+        # output = self.output(h)
+        return h

gpst/backend_loader.py

@@ -0,0 +1,42 @@
+import importlib
+import pathlib
+import sys
+import subprocess
+
+# gpst/backend_loader.py
+import os
+import glob
+import torch
+from torch.utils.cpp_extension import load
+from setuptools import setup
+from torch.utils.cpp_extension import BuildExtension, CUDAExtension, CppExtension
+
+def load_cpp_backend2():
+    include_dirs = os.path.dirname(os.path.abspath(__file__))
+    source_files = glob.glob(os.path.join(include_dirs, "cpp_extension", "*.cpp"))
+
+    if not source_files:
+        raise FileNotFoundError("No C++ source files found in cpp_extension/")
+
+    print("[INFO] Building cppbackend...")
+    cppbackend = load(
+        name="cppbackend",
+        sources=source_files,
+        extra_include_paths=[include_dirs],  # same as include_dirs in setup.py
+        verbose=True
+    )
+    return cppbackend
+
+
+def load_cpp_backend():
+    try:
+        # Try importing the compiled module
+        import gpst.cppbackend
+        #return importlib.import_module("cppbackend")
+    except ModuleNotFoundError:
+        print("[INFO] C++ backend not found. Building now...")
+        repo_root = pathlib.Path(__file__).parent.resolve()
+        setup_script = repo_root / "setup.py"
+        subprocess.check_call([sys.executable, str(setup_script), "build_ext", "--inplace"])
+
+        import gpst.cppbackend

gpst/config.py

@@ -0,0 +1,10 @@
+from transformers import PretrainedConfig
+
+class GPSTConfig(PretrainedConfig):
+    model_type = "gpst"
+    
+    def __init__(self, r2d2=None, gpt=None, **kwargs):#, gptconfig, r2d2config, **kwargs):
+
+        self.gptconfig = gpt
+        self.r2d2config = r2d2
+        super().__init__(**kwargs)

gpst/cpp_extension/binding.cpp

@@ -0,0 +1,29 @@
+#include <torch/torch.h>
+#include "py_backend.h"
+#include <pybind11/pybind11.h>
+#include <pybind11/numpy.h>
+#include <pybind11/stl.h>
+
+
+PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
+    std::string name = std::string("TableManager");
+    py::class_<TableManager>(m, name.c_str())
+        .def(py::init([](const py::array_t<int>& seq_lens, const py::array_t<int>& group_ids,
+                         const py::array_t<int>& merge_orders, 
+                         size_t window_size, size_t cache_id_offset, size_t detach_cache_id_offset,
+                         vector<py::array_t<int>>& span_ids)
+                      { return new TableManager(seq_lens, group_ids, merge_orders, window_size, 
+                                                cache_id_offset, detach_cache_id_offset, span_ids); }))
+        .def("step", &TableManager::step)
+        .def("root_ids", &TableManager::root_ids)
+        .def("is_finished", &TableManager::is_finished)
+        .def("prepare_bilm", &TableManager::prepare_bilm)
+        .def("prepare_generation", &TableManager::prepare_generation)
+        .def("batch_size", &TableManager::batch_size);
+    name = std::string("SpanTokenizer");
+    py::class_<SpanTokenizer>(m, name.c_str())
+        .def(py::init([](vector<py::array_t<int>>& dictionary, int max_entry_id) {
+            return new SpanTokenizer(dictionary, max_entry_id);
+        }))
+        .def("tokenize", &SpanTokenizer::tokenize);
+}

gpst/cpp_extension/py_backend.cpp

@@ -0,0 +1,1100 @@
+// Copyright (c) 2024 Ant Group
+// Author: Xiang Hu
+// #undef NDEBUG
+#include <stdexcept>
+#include <cstring>
+#include <numeric>
+#include "py_backend.h"
+#include <cassert>
+#include <cmath>
+#include <functional>
+using namespace torch::indexing;
+
+#ifndef NDEBUG
+#   define ASSERT(condition, message) \
+    do { \
+        if (! (condition)) { \
+            std::cerr << "Assertion `" #condition "` failed in " << __FILE__ \
+                      << " line " << __LINE__ << ": " << message << std::endl; \
+            std::terminate(); \
+        } \
+    } while (false)
+#else
+#   define ASSERT(condition, message) do { } while (false)
+#endif
+
+LinkedNode::LinkedNode(Cell * value):m_pLeft{NULL}, m_pRight{NULL}, m_pLeftup{NULL}, 
+m_pRightup{NULL}, m_pLeftdown{NULL}, m_pRightdown{NULL}, m_pCell{value} {
+}
+
+LinkedNode::~LinkedNode(){
+    this->m_pCell->setNode(NULL);
+}
+
+Cell * LinkedNode::getCell() const {
+    return this->m_pCell;
+}
+
+LinkedNode * LinkedNode::left() const {
+    return this->m_pLeft;
+}
+
+LinkedNode * LinkedNode::right() const {
+    return this->m_pRight;
+}
+
+LinkedNode * LinkedNode::leftup() const {
+    return this->m_pLeftup;
+}
+
+LinkedNode * LinkedNode::rightup() const {
+    return this->m_pRightup;
+}
+
+LinkedNode * LinkedNode::leftdown() const {
+    return this->m_pLeftdown;
+}
+
+LinkedNode * LinkedNode::rightdown() const {
+    return this->m_pRightdown;
+}
+
+void LinkedNode::setLeft(LinkedNode * other) {
+    if (this->m_pLeft != other) {
+        this->m_pLeft = other;
+        if (other != NULL) {
+            other->setRight(this);
+        }
+    }
+}
+
+void LinkedNode::setRight(LinkedNode * other) {
+    if (this->m_pRight != other) {
+        this->m_pRight = other;
+        if (other != NULL) {
+            other->setLeft(this);
+        }
+    }
+}
+
+void LinkedNode::setLeftup(LinkedNode * other) {
+    if (this->m_pLeftup != other) {
+        this->m_pLeftup = other;
+        if (other != NULL) {
+            assert(this->getCell()->j == other->getCell()->j);
+            other->setRightdown(this);
+        }
+    }
+}
+
+void LinkedNode::setRightup(LinkedNode * other) {
+    if (this->m_pRightup != other) {
+        this->m_pRightup = other;
+        if (other != NULL) {
+            assert(this->getCell()->i == other->getCell()->i);
+            other->setLeftdown(this);
+        }
+    }
+}
+
+void LinkedNode::setLeftdown(LinkedNode * other) {
+    if (this->m_pLeftdown != other) {
+        this->m_pLeftdown = other;
+        if (other != NULL) {
+            assert(this->getCell()->i == other->getCell()->i);
+            other->setRightup(this);
+        }
+    }
+}
+
+void LinkedNode::setRightdown(LinkedNode * other) {
+    if (this->m_pRightdown != other) {
+        this->m_pRightdown = other;        
+        if (other != NULL) {
+            assert(this->getCell()->j == other->getCell()->j);
+            other->setLeftup(this);
+        }
+    }
+}
+
+Cell::Cell(int i, int j, int window_size, TableManager * mgr, int batch_id, bool is_root):i{i}, j{j}, detached{false}, 
+best_split{-1}, a_ij_split{-1}, cache_id{-1}, batch_id(batch_id), m_pReadyChild(0), m_pMgr(mgr), m_bIsRoot(is_root), ext_vocab_id(0) {
+    this->m_pNode = NULL;
+    if (j > i) {
+        this->split_size = j - i < window_size ? j - i : window_size;
+        this->splits = new int[this->split_size];
+        memset(this->splits, 0, this->split_size * sizeof(int));
+    } else {
+        this->split_size = 0;
+        this->splits = 0;
+    }
+}
+
+Cell::~Cell() {
+    if (this->splits != NULL) {
+        delete this->splits;
+    }
+    if (this->m_pNode != NULL) {
+        delete this->m_pNode;
+    }
+}
+
+int Cell::getDetachedCacheID(int detach_offset) const {
+    if (this->detached || this->i == this->j) {
+        return detach_offset + this->cache_id;
+    } else {
+        return this->cache_id;
+    }
+}
+
+LinkedNode * Cell::getNode() const {
+    return this->m_pNode;
+}
+
+void Cell::setNode(LinkedNode * target) {
+    assert(this->m_pNode == NULL || target == NULL);
+    assert(this->m_pNode != NULL || target != NULL);
+    this->m_pNode = target;
+}
+
+void Cell::addParent(Cell * parent) {
+    this->m_lParents.push_back(parent);
+}
+
+void Cell::onReady() {
+    for (auto const & cell : this->m_lParents) {
+        cell->notifyChildReady();
+    }
+}
+
+void Cell::notifyChildReady() {
+    ++this->m_pReadyChild;
+    assert(this->m_pReadyChild <= 2 * this->split_size);
+    if (this->m_pReadyChild == 2 * this->split_size && (m_bIsRoot || this->m_lParents.size() > 0)) {
+        // not root and has parents
+        this->m_pMgr->on_cell_ready(this);
+    }
+}
+
+
+int Cell::getBestSplit() const {
+    int k = this->best_split;
+    assert(k >= 0);
+    assert(k < this->split_size);
+    return this->splits[k];
+}
+
+int Cell::getGumbelSplit() const {
+    int k = this->a_ij_split;
+    assert(k >= 0);
+    assert(k < this->split_size);
+    return this->splits[k];
+}
+
+CellTable::CellTable(int seq_len, int window_size, int batch_i, TableManager * mgr):m_iCellOffset{0}, m_iMaxCreatedCells{2 * (window_size + 1) * seq_len}, 
+m_iSeqLen{seq_len}, m_iWindowSize{window_size}, m_iBatchId(batch_i), m_pMgr(mgr) {
+    this->m_pCells = new Cell*[seq_len * seq_len];
+    memset(this->m_pCells, 0, seq_len * seq_len * sizeof(Cell*));
+
+    this->m_pCreatedCells = new Cell*[this->m_iMaxCreatedCells];
+    memset(this->m_pCreatedCells, 0, this->m_iMaxCreatedCells * sizeof(Cell*));
+}
+
+CellTable::~CellTable() {
+    for (int i = 0; i < this->m_iCellOffset; ++i) {
+        assert(this->m_pCreatedCells[i] != NULL);
+        delete this->m_pCreatedCells[i];
+    }
+    delete this->m_pCreatedCells;
+    delete this->m_pCells;
+}
+
+Cell * CellTable::get(const int i, const int j) {
+    assert(i <= j);
+    assert(j < this->m_iSeqLen);
+    if (this->m_pCells[i * this->m_iSeqLen + j] == NULL) {
+        bool is_root = j - i + 1 == this->m_iSeqLen && i == 0;
+        Cell * new_cell = new Cell(i, j, this->m_iWindowSize, this->m_pMgr, this->m_iBatchId, is_root);
+        this->m_pCreatedCells[this->m_iCellOffset++] = new_cell;
+        this->m_pCells[i * this->m_iSeqLen + j] = new_cell;
+    }
+    return this->m_pCells[i * this->m_iSeqLen + j];
+}
+
+bool CellTable::isEmpty(const int i, const int j) {
+    return this->m_pCells[i * this->m_iSeqLen + j] == NULL;
+}
+
+int CellTable::getLen() const {
+    return this->m_iSeqLen;
+}
+
+void init_active_cells(int window_size, int seq_len, CellTable * cell_table) {
+    for (int layer_i = 0; layer_i <= window_size; ++layer_i) {
+        LinkedNode * left_previous = NULL;
+        for (int pos_i = 0; pos_i < seq_len - layer_i; ++pos_i) {
+            Cell * cell_ij = cell_table->get(pos_i, pos_i + layer_i);
+            if (layer_i > 0) {
+                for (int split_idx = 0; split_idx < layer_i; ++split_idx) {
+                    cell_ij->splits[split_idx] = pos_i + split_idx;
+                }
+            }
+            LinkedNode * node = new LinkedNode(cell_ij);
+            cell_ij->setNode(node);
+            node->setLeft(left_previous);
+            left_previous = node;
+            LinkedNode * leftdown = NULL;
+            LinkedNode * rightdown = NULL;
+            if (layer_i > 0) {
+                leftdown = cell_table->get(pos_i, pos_i + layer_i - 1)->getNode();
+                rightdown = cell_table->get(pos_i + 1, pos_i + layer_i)->getNode();
+            }
+            node->setLeftdown(leftdown);
+            node->setRightdown(rightdown);
+        }
+    }
+}
+
+LinkedNode * create_new_node(LinkedNode * leftdown, LinkedNode * rightdown, LinkedNode * left, 
+LinkedNode * right, LinkedNode * ld_most, LinkedNode * rd_most, CellTable * table) {
+    int i = leftdown->getCell()->i;
+    int j = rightdown->getCell()->j;
+    Cell * new_cell = table->get(i, j);
+    LinkedNode * new_node = new LinkedNode(new_cell);
+    LinkedNode * current = ld_most;
+
+    int idx = 0;
+    while (current != rd_most->right()) {
+        assert(idx < new_cell->split_size);
+        new_cell->splits[idx++] = current->getCell()->j;
+        current = current->right();
+    }
+
+    new_cell->setNode(new_node);
+    new_node->setLeftdown(leftdown);
+    new_node->setRightdown(rightdown);
+    new_node->setLeft(left);
+    new_node->setRight(right);
+    return new_node;
+}
+
+void prune(LinkedNode * node, std::function<void(Cell*)> on_new_cell, CellTable * table) {
+    node->getCell()->detached = true;
+    LinkedNode * leftdown = node->leftdown();
+    LinkedNode * rightdown = node->rightdown();
+
+    LinkedNode * left_node = leftdown;
+    LinkedNode * tmp = NULL;
+
+    int left_steps = 0;
+    while (left_node != NULL) {
+        if (left_node->left() != NULL) {
+            left_node->left()->setRight(left_node->rightup());
+        }
+        if (left_node->rightup() != NULL) {
+            left_node->rightup()->setLeftdown(left_node->leftdown());
+        }
+        left_steps += 1;
+
+        if (left_node->leftup() != NULL) {
+            tmp = left_node->leftup();
+            delete left_node;
+            left_node = tmp;
+        } else {
+            if (left_node->rightup() != NULL) {
+                tmp = left_node->rightup();
+                left_steps += 1;
+            } else {
+                tmp = left_node->right();
+            }
+            // std::cout << "delete: " << left_node->getCell()->i << "," << left_node->getCell()->j << std::endl;
+            delete left_node;
+            left_node = tmp;
+            break;
+        }
+    }
+
+    // std::cout << "left_node: " << left_node->getCell()->i << "," << left_node->getCell()->j << std::endl;
+
+    while (left_node->rightup() != NULL) {
+        left_node = left_node->rightup();
+        // std::cout << "left_node: " << left_node->getCell()->i << "," << left_node->getCell()->j << std::endl;
+        left_steps += 1;
+    }
+    // std::cout << "pb" << std::endl;
+
+    LinkedNode * right_node = rightdown;
+    assert(rightdown != NULL);
+    int right_steps = 0;
+    while (right_node != NULL) {
+        // std::cout << "right_node: " << right_node->getCell()->i << "," << right_node->getCell()->j << std::endl;
+        if (right_node->right() != NULL) {
+            right_node->right()->setLeft(right_node->leftup());
+        }
+        if (right_node->leftup() != NULL) {
+            right_node->leftup()->setRightdown(right_node->rightdown());
+        }
+        right_steps += 1;
+
+        if (right_node->rightup() != NULL) {
+            tmp = right_node->rightup();
+            delete right_node;
+            right_node = tmp;
+        } else {
+            if (right_node->leftup() != NULL) {
+                tmp = right_node->leftup();
+                right_steps += 1;
+            } else {
+                tmp = right_node->left();
+            }
+            // std::cout << "delete : " << right_node->getCell()->i << "," << right_node->getCell()->j << std::endl;
+            delete right_node;
+            right_node = tmp;
+            break;
+        }
+    }
+    // std::cout << "right_node: " << right_node->getCell()->i << "," << right_node->getCell()->j << std::endl;
+    while (right_node->leftup() != NULL) {
+        // std::cout << "right_node: " << right_node->getCell()->i << "," << right_node->getCell()->j << std::endl;
+        right_node = right_node->leftup();
+        right_steps += 1;
+    }
+
+    // std::cout << "pc" << std::endl;
+
+    LinkedNode * current = left_node->left() != NULL ? left_node->left() : left_node;
+    LinkedNode * end = right_node->right() != NULL ? right_node->right() : right_node;
+
+    LinkedNode * current_ld_most = current;
+    while (current_ld_most->leftdown() != NULL) {
+        current_ld_most = current_ld_most->leftdown();
+    }
+
+    LinkedNode * current_rd_most = current;
+    while (current_rd_most->rightdown() != NULL) {
+        current_rd_most = current_rd_most->rightdown();
+    }
+
+    // std::cout << "pd" << std::endl;
+    while (current != end) {
+        LinkedNode * node_left = current->leftup();
+        LinkedNode * node_right = current->right()->rightup();
+        // std::cout << "pcreate" << std::endl;
+        LinkedNode * node = create_new_node(current, current->right(), node_left, node_right,
+                                            current_ld_most, current_rd_most, table);
+        // std::cout << "pcreate over" << std::endl;
+        on_new_cell(node->getCell());
+        current = current->right();
+        current_ld_most = current_ld_most->right();
+        current_rd_most = current_rd_most->right();
+
+    }
+}
+
+
+TableManager::TableManager(const py::array_t<int>& seq_lens, const py::array_t<int>& group_ids, const py::array_t<int>& merge_orders, const int window_size, 
+const int cache_id_offset, const int detach_id_offset, vector<py::array_t<int>>& span_ids): m_iBatchSize{seq_lens.shape()[0]}, m_iWindowSize{window_size}, 
+m_iCacheOffset{cache_id_offset}, m_iCurrentStep{1}, m_iDetachCacheOffset(detach_id_offset), m_iCellNum{0} {
+    this->m_pCellTables = new CellTable*[this->m_iBatchSize];
+    // this->m_pMergeOrders = new Span*[this->m_iBatchSize];
+    auto pMergeOrders = new Span*[this->m_iBatchSize];
+
+    auto buf = seq_lens.request();
+    int * seq_lens_ptr = (int*)buf.ptr;
+    int max_seq_len = 0;
+    int seq_len_sum = 0;
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        int seq_len = seq_lens_ptr[batch_i];
+        CellTable * table = new CellTable(seq_len, window_size, batch_i, this);
+        this->m_pCellTables[batch_i] = table;
+        init_active_cells(window_size, seq_len, table);
+        for (int pos = 0; pos < seq_len; ++pos) {
+            table->get(pos, pos)->cache_id = this->m_iCacheOffset + this->m_iCellNum++;
+        }
+
+        if (span_ids.size() > 0) {
+            buf = span_ids[batch_i].request();
+            int total_size = span_ids[batch_i].shape(0) / 3;
+            // std::cout << "batch: " << batch_i << " total size: " << total_size << std::endl;
+            // std::cout << "table seq len: " << table->getLen() << std::endl;
+            int * span_ids_ptr = (int*)buf.ptr;
+            for (int span_idx = 0; span_idx < total_size; ++span_idx) {
+                // std::cout << span_ids_ptr[span_idx * 3] << ", " << span_ids_ptr[span_idx * 3 + 1] << "->" << span_ids_ptr[span_idx * 3 + 2] << ";";
+                if (!table->isEmpty(span_ids_ptr[span_idx * 3], span_ids_ptr[span_idx * 3 + 1])) {
+                    table->get(span_ids_ptr[span_idx * 3], span_ids_ptr[span_idx * 3 + 1])->ext_vocab_id = span_ids_ptr[span_idx * 3 + 2];
+                }
+            }
+        }
+
+        max_seq_len = seq_len > max_seq_len ? seq_len : max_seq_len;
+        seq_len_sum += seq_len;
+    }
+
+    buf = group_ids.request();
+    int * merge_order_ptr = 0;
+    int merge_order_L = 0;
+    int * group_ids_ptr = (int*)buf.ptr;
+    int max_group_size = group_ids_ptr[m_iBatchSize - 1] + 1;
+
+    if (m_iBatchSize != max_group_size) {
+        merge_order_L = max_seq_len - 1;
+        merge_order_ptr = new int[m_iBatchSize * merge_order_L];
+        std::fill_n(merge_order_ptr, m_iBatchSize * merge_order_L, -1);
+        buf = merge_orders.request();
+        int * chunk_merge_order_ptr = (int*)buf.ptr;
+        int chunk_size = merge_orders.shape(1); // sum(seq_lens) - 1
+
+        int prev_group_id = -1;
+        int offset = 0;
+        int * pos_chunk2sent = new int[max_group_size * chunk_size];
+        int * bid_chunk2sent = new int[max_group_size * chunk_size];
+        std::fill_n(bid_chunk2sent, max_group_size * chunk_size, -1);
+        std::fill_n(pos_chunk2sent, max_group_size * chunk_size, -1);
+        for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+            if (group_ids_ptr[batch_i] != prev_group_id) {
+                prev_group_id = group_ids_ptr[batch_i];
+                offset = 0;
+            }
+
+            for (int sent_idx = 0; sent_idx < seq_lens_ptr[batch_i] - 1; ++sent_idx) {
+                if (offset + sent_idx < chunk_size) {
+                    pos_chunk2sent[prev_group_id * chunk_size + offset + sent_idx] = sent_idx;
+                    bid_chunk2sent[prev_group_id * chunk_size + offset + sent_idx] = batch_i;   
+                }
+            }
+            offset += seq_lens_ptr[batch_i];
+        }
+
+        int * sent_offset = new int[m_iBatchSize];
+        std::fill_n(sent_offset, m_iBatchSize, 0);
+        for (int group_i = 0; group_i < max_group_size; ++group_i) {
+            for (int pos = 0; pos < chunk_size; ++pos) {
+                assert(chunk_merge_order_ptr[group_i * chunk_size + pos] >= 0 && chunk_merge_order_ptr[group_i * chunk_size + pos] < chunk_size);
+                int sent_id = bid_chunk2sent[group_i * chunk_size + chunk_merge_order_ptr[group_i * chunk_size + pos]];
+                int split_pos = pos_chunk2sent[group_i * chunk_size + chunk_merge_order_ptr[group_i * chunk_size + pos]];
+                // std::cout << "sent_id" << sent_id << ", " << split_pos << std::endl;
+                if (sent_id != -1 && split_pos < seq_lens_ptr[sent_id] - 1) {
+                    // std::cout << sent_id << "|" << sent_offset[sent_id] << "|" << split_pos << ", " << std::endl;
+                    assert(sent_offset[sent_id] < merge_order_L);
+                    assert(sent_id >= 0 && sent_id < m_iBatchSize);
+                    merge_order_ptr[sent_id * merge_order_L + sent_offset[sent_id]++] = split_pos;
+                }
+            }
+        }
+        delete pos_chunk2sent;
+        delete bid_chunk2sent;
+        delete sent_offset;
+    } else {
+        buf = merge_orders.request();
+        merge_order_ptr = (int*)buf.ptr;
+        merge_order_L = merge_orders.shape()[1];
+    }
+    // std::cout << "B" << std::endl;
+
+    //convert merge order to cell i,j
+    int left_i = 0;
+    int right_j = 0;
+    int merge_pos = 0;
+    // assert(seq_lens.shape()[0] == merge_orders.shape()[0]);
+    
+    
+    int * current_merge_orders = 0;
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        int seq_len = seq_lens_ptr[batch_i];
+        Span * left_splits = new Span[seq_len - 1];
+        Span * right_splits = new Span[seq_len - 1];
+        current_merge_orders = merge_order_ptr + batch_i * merge_order_L;
+        for (int split = 0; split < seq_len - 1; ++split) {
+            left_splits[split].i = split;
+            left_splits[split].j = split;
+            right_splits[split].i = split + 1;
+            right_splits[split].j = split + 1;
+        }
+
+        Span * merge_orders = new Span[seq_len - 1];
+        for (int action_i = 0; action_i < seq_len - 1; ++action_i) {
+            merge_pos = current_merge_orders[action_i];
+            assert(merge_pos < seq_len - 1);
+            left_i = left_splits[merge_pos].i;
+            right_j = right_splits[merge_pos].j;
+
+            merge_orders[action_i].i = left_i;
+            merge_orders[action_i].j = right_j;
+            if (left_i >= 1) {
+                right_splits[left_i - 1].i = left_i;
+                right_splits[left_i - 1].j = right_j;
+            }
+            if (right_j < seq_len - 1) {
+                left_splits[right_j].i = left_i;
+                left_splits[right_j].j = right_j; 
+            }
+            // std::cout << "merge span: " << left_i << ", " << right_j << std::endl; 
+        }
+        pMergeOrders[batch_i] = merge_orders;
+
+        delete left_splits;
+        delete right_splits;
+    }
+    // std::cout << "C" << std::endl;
+
+    // int max_seq_len = *max_element(seq_lens.begin(), seq_lens.end());
+    // int seq_len_sum = accumulate(seq_lens.begin(), seq_lens.end(), 0);
+    this->m_iMaxSeqLen = max_seq_len;
+    this->m_pCellOrders = new Cell**[max_seq_len - 1];
+    this->m_pCellNums = new int[max_seq_len - 1];
+    this->m_pTargetCacheIds = new long*[max_seq_len - 1];
+    this->m_pGroupCacheIds = new long*[max_seq_len - 1];
+    this->m_pTargetExtIds = new long*[max_seq_len - 1];
+    this->m_pDetachGroupCacheIds = new long*[max_seq_len - 1];
+
+    this->m_pLDRCache_ids = 0;
+    this->m_pExtIds = 0;
+    this->m_pPositionIds = 0;
+    this->m_pTgtIds = 0;
+
+    this->m_pSpanMasks = 0;
+    this->m_pSplitTargets = 0;
+    this->m_pSpanGatherIds = 0;
+    this->m_pTokenPositions = 0;
+
+    for (int step = 1; step < max_seq_len; ++step) {
+        this->m_pCellOrders[step - 1] = 0;
+        this->m_pCellNums[step - 1] = 0;
+        this->m_pTargetCacheIds[step - 1] = 0;
+        this->m_pGroupCacheIds[step - 1] = 0;
+        this->m_pTargetExtIds[step - 1] = 0;
+        this->m_pDetachGroupCacheIds[step - 1] = 0;
+    }
+
+    this->build_cell_dependencies(pMergeOrders);
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        int seq_len = seq_lens_ptr[batch_i];
+        CellTable * table = this->m_pCellTables[batch_i];
+        for (int pos = 0; pos < seq_len; ++pos) {
+            table->get(pos, pos)->onReady();
+        }
+    }
+
+    if (m_iBatchSize != max_group_size) {
+        delete merge_order_ptr;
+    }
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        delete pMergeOrders[batch_i];
+    }
+    delete pMergeOrders;
+}
+
+TableManager::~TableManager() {
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        delete this->m_pCellTables[batch_i];
+        // delete this->m_pMergeOrders[batch_i];
+    }
+    for (int step = 0; step < this->m_iMaxSeqLen - 1; ++step) {
+        if (this->m_pCellOrders[step] != 0)
+            delete this->m_pCellOrders[step];
+        if (this->m_pTargetCacheIds[step] != 0)
+            delete this->m_pTargetCacheIds[step];
+        if (this->m_pGroupCacheIds[step] != 0)
+            delete this->m_pGroupCacheIds[step];
+        if (this->m_pTargetExtIds[step] != 0)
+            delete this->m_pTargetExtIds[step];
+        if (this->m_pDetachGroupCacheIds[step] != 0)
+            delete this->m_pDetachGroupCacheIds[step];
+    }
+
+
+    if (m_pLDRCache_ids != 0)
+        delete m_pLDRCache_ids;
+    if (m_pExtIds != 0)
+        delete m_pExtIds;
+    if (m_pPositionIds != 0)
+        delete m_pPositionIds;
+    if (m_pTgtIds != 0)
+        delete m_pTgtIds;
+    if (m_pSpanMasks != 0)
+        delete m_pSpanMasks;
+    if (m_pSplitTargets != 0)
+        delete m_pSplitTargets;
+    if (m_pSpanGatherIds != 0)
+        delete m_pSpanGatherIds;
+    if (m_pTokenPositions != 0)
+        delete m_pTokenPositions;
+
+    delete this->m_pCellNums;
+    delete this->m_pCellOrders;
+    delete this->m_pCellTables;
+    // delete this->m_pMergeOrders;
+    delete this->m_pTargetCacheIds;
+    delete this->m_pTargetExtIds;
+    delete this->m_pGroupCacheIds;
+    delete this->m_pDetachGroupCacheIds;
+}
+
+void TableManager::push_cell(Cell * cell) {
+    int current_cache_id = this->m_iCacheOffset + this->m_iCellNum++;
+    cell->cache_id = current_cache_id;
+}
+
+void TableManager::build_cell_dependencies(Span ** pMergeOrders) {
+    for (int step = 1; step < this->m_iMaxSeqLen; ++step) {
+        // #pragma omp parallel for
+        for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+            CellTable * table = this->m_pCellTables[batch_i];
+            int seq_len = table->getLen();
+            // printf("batch_i: %d, seq_len: %d\n", batch_i, seq_len);
+            if (step < seq_len) {
+                if (step <= this->m_iWindowSize) {
+                    for (int i = 0; i < seq_len - step; ++i) {
+                        int j = i + step;
+                        assert(j < seq_len);
+                        Cell * current_cell = table->get(i, j);
+                        for (int split_idx = 0; split_idx < current_cell->split_size; ++split_idx) {
+                            int k = current_cell->splits[split_idx];
+                            assert(i <= k);
+                            assert(k < j);
+                            Cell * cell_ik = table->get(i, k);
+                            Cell * cell_kj = table->get(k + 1, j);
+                            cell_ik->addParent(current_cell);
+                            cell_kj->addParent(current_cell);
+                        }
+                    }
+                } else {
+                    Span & merge_span = pMergeOrders[batch_i][step - this->m_iWindowSize - 1];
+                    int i = merge_span.i;
+                    int j = merge_span.j;
+                    // printf("(%d, %d)\n", i, j);
+                    prune(table->get(i, j)->getNode(), [&](Cell * cell) {
+                        for (int sp_idx = 0; sp_idx < cell->split_size; ++sp_idx) {
+                            int k = cell->splits[sp_idx];
+                            assert(cell->i <= k);
+                            assert(k < cell->j);
+                            Cell * cell_ik = table->get(cell->i, k);
+                            Cell * cell_kj = table->get(k + 1, cell->j);
+                            cell_ik->addParent(cell);
+                            cell_kj->addParent(cell);
+                        }
+                        
+                    }, table);
+                }
+            }
+        }
+    }
+}
+
+void TableManager::on_cell_ready(Cell * ready_cell) {
+    this->m_lReadyCells.push_back(ready_cell);
+}
+
+bool TableManager::is_finished() {
+    return this->m_lReadyCells.size() == 0;
+}
+
+vector<at::Tensor> TableManager::step() {
+    int current_step = this->m_iCurrentStep;
+
+    int total_size = this->m_lReadyCells.size();
+    int group_size = this->m_iCurrentStep <= this->m_iWindowSize ? this->m_iCurrentStep : this->m_iWindowSize;
+    this->m_pCellNums[current_step - 1] = total_size;
+    this->m_pCellOrders[current_step - 1] = new Cell*[total_size];
+
+    // at::Tensor target_cache_ids_ = torch::zeros({total_size}, at::kLong);
+    // at::Tensor group_cache_ids_ = torch::zeros({total_size, group_size, 2}, at::kLong);
+    // at::Tensor detach_group_cache_ids_ = torch::zeros({total_size, group_size, 2}, at::kLong);
+
+    this->m_pTargetCacheIds[current_step - 1] = new long[total_size];
+    this->m_pGroupCacheIds[current_step - 1] = new long[total_size * group_size * 2];
+    this->m_pTargetExtIds[current_step - 1] = new long[total_size];
+    this->m_pDetachGroupCacheIds[current_step - 1] = new long[total_size * group_size * 2];
+
+    auto tgt_cache_ids_ptr = this->m_pTargetCacheIds[current_step - 1];
+    auto tgt_ext_id_ptr = this->m_pTargetExtIds[current_step - 1];
+    auto group_ids_ptr = this->m_pGroupCacheIds[current_step - 1];
+    auto detach_group_ids_ptr = this->m_pDetachGroupCacheIds[current_step - 1];
+
+    int idx_offset = 0;
+    for (int cell_idx = 0; cell_idx < total_size; ++cell_idx) {
+        Cell * cell_ptr = this->m_lReadyCells.front();
+        this->m_lReadyCells.pop_front();
+        CellTable * table = this->m_pCellTables[cell_ptr->batch_id];
+        int i = cell_ptr->i;
+        int j = cell_ptr->j;
+        cell_ptr->onReady();
+        this->push_cell(cell_ptr);
+        this->m_pCellOrders[this->m_iCurrentStep - 1][cell_idx] = cell_ptr;
+        tgt_cache_ids_ptr[cell_idx] = cell_ptr->cache_id;
+        tgt_ext_id_ptr[cell_idx] = cell_ptr->ext_vocab_id;
+        for (int split_idx = 0; split_idx < cell_ptr->split_size; ++split_idx) {
+            int k = cell_ptr->splits[split_idx];
+            assert(i <= k);
+            assert(k < j);
+            Cell * cell_ik = table->get(i, k);
+            Cell * cell_kj = table->get(k + 1, j);
+            // group_cache_ids_.index({cell_idx, split_idx, 0}) = cell_ik->cache_id;
+            // group_cache_ids_.index({cell_idx, split_idx, 1}) = cell_kj->cache_id;
+            // detach_group_cache_ids_.index({cell_idx, split_idx, 0}) = cell_ik->getDetachedCacheID(this->m_iDetachCacheOffset);
+            // detach_group_cache_ids_.index({cell_idx, split_idx, 1}) = cell_kj->getDetachedCacheID(this->m_iDetachCacheOffset);
+            group_ids_ptr[idx_offset] = cell_ik->cache_id;
+            group_ids_ptr[idx_offset + 1] = cell_kj->cache_id;
+            detach_group_ids_ptr[idx_offset] = cell_ik->getDetachedCacheID(this->m_iDetachCacheOffset);
+            detach_group_ids_ptr[idx_offset + 1] = cell_kj->getDetachedCacheID(this->m_iDetachCacheOffset);
+
+            idx_offset += 2;
+        }
+    }
+
+    at::Tensor target_cache_ids_ = torch::from_blob(tgt_cache_ids_ptr, {total_size}, at::kLong);
+    at::Tensor group_cache_ids_ = torch::from_blob(group_ids_ptr, {total_size, group_size, 2}, at::kLong);
+    at::Tensor ext_id_cache_ids_ = torch::from_blob(tgt_ext_id_ptr, {total_size}, at::kLong);
+    at::Tensor detach_group_cache_ids_ = torch::from_blob(detach_group_ids_ptr, {total_size, group_size, 2}, at::kLong);
+
+    this->m_iCurrentStep += 1;
+    return {target_cache_ids_, ext_id_cache_ids_, group_cache_ids_, detach_group_cache_ids_};
+    // return {target_cache_ids_, group_cache_ids_, detach_group_cache_ids_};
+}
+
+bool hit_span(int i, int k, int j, py::array_t<int>& atom_spans) {
+    auto atom_span_pt = atom_spans.unchecked<2>();
+    int atom_span_st = 0, atom_span_ed = 0;
+    for (int atom_i = 0; atom_i < atom_spans.shape(0); ++atom_i) {
+        atom_span_st = atom_span_pt(atom_i, 0);
+        atom_span_ed = atom_span_pt(atom_i, 1);
+
+        if (j < atom_span_st || i > atom_span_ed || (i >= atom_span_st && j <= atom_span_ed)) {
+            // no overlap
+            continue;
+        }
+        if ((k < atom_span_st && j >= atom_span_ed) || (k + 1 > atom_span_ed && i <= atom_span_st)) {
+            continue;
+        }
+        return true;
+    }
+    return false;
+}
+
+void induce_best_splits(int max_seq_len, int * cell_nums, Cell *** cell_orders, vector<py::array_t<int>> &best_splits, 
+    vector<py::array_t<int>> &a_ij_splits, vector<py::array_t<int>> &atom_spans) {
+    // a_ij after gumbel_softmax maybe different from best_splits
+    for (int step = 0; step < max_seq_len - 1; ++step) {
+        for (int cell_idx = 0; cell_idx < cell_nums[step]; ++cell_idx) {
+            Cell * current_cell = cell_orders[step][cell_idx];
+            assert(current_cell->i <= current_cell->j);
+            // atom_spans[current_cell->batch_id]
+            auto best_splits_mat = best_splits[step].unchecked<2>();
+            auto a_ij_np = a_ij_splits[step].unchecked<2>();
+            current_cell->a_ij_split = a_ij_np(cell_idx, 0);
+            if (atom_spans.size() > 0 && atom_spans[current_cell->batch_id].shape(0) > 0) {
+                for (int split_idx = 0; split_idx < best_splits[step].shape(1); ++split_idx) {
+                    int split = best_splits_mat(cell_idx, split_idx);
+                    int k = current_cell->splits[split];
+                    if (!hit_span(current_cell->i, k, current_cell->j, atom_spans[current_cell->batch_id])) {
+                        current_cell->best_split = split;
+                        break;
+                    }
+                }
+                // assert(current_cell->best_split != -1);
+            } else {
+                current_cell->best_split = best_splits_mat(cell_idx, 0);
+            }
+        }
+    }
+}
+
+at::Tensor TableManager::root_ids() {
+    at::Tensor t = torch::zeros(this->m_iBatchSize, torch::kLong);
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        CellTable * tbl = this->m_pCellTables[batch_i];
+        t[batch_i] = tbl->get(0, tbl->getLen() - 1)->cache_id;
+    }
+    return t;
+}
+
+const int TableManager::batch_size() const {
+    return this->m_iBatchSize;
+}
+
+int left_most(CellTable * table, int idx, int bos_id, int eos_id) {
+    if (idx < 0) {
+        return bos_id;
+    }
+    assert (idx < table->getLen() - 1);
+    for (int start = 0; start <= idx; ++start) {
+        if (!table->isEmpty(start, idx)) {
+            return table->get(start, idx)->cache_id;
+        }
+    }
+    assert (false);
+}
+
+int right_most(CellTable * table, int idx, int bos_id, int eos_id) {
+    if (idx >= table->getLen()) {
+        return eos_id;
+    }
+    assert (idx > 0);
+    for (int end = table->getLen() - 1; end >= idx; --end) {
+        if (!table->isEmpty(idx, end)) {
+            return table->get(idx, end)->cache_id;
+        }
+    }
+    assert (false);
+}
+
+at::Tensor TableManager::prepare_bilm(int total_len, int bos_id, int eos_id) {
+    at::Tensor cache_ids = torch::zeros({total_len, 2}, torch::kLong);
+    int offset = 0;
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        CellTable * tbl = this->m_pCellTables[batch_i];
+        for (int idx = 0; idx < tbl->getLen(); ++idx) {
+            cache_ids.index({offset, 0}) = left_most(tbl, idx - 1, bos_id, eos_id);
+            cache_ids.index({offset, 1}) = right_most(tbl, idx + 1, bos_id, eos_id);
+            ++offset;
+        }
+    }
+    assert(offset == total_len);
+    return cache_ids;
+}
+
+vector<at::Tensor> TableManager::prepare_generation(vector<py::array_t<int>>& score_splits, 
+                                                    vector<py::array_t<int>>& a_ij_splits, 
+                                                    vector<py::array_t<int>>& atom_spans, 
+                                                    const py::array_t<int>& input_ids,
+                                                    const py::array_t<int>& group_ids,
+                                                    const py::array_t<int>& eos_labels,
+                                                    const int reduce_id,
+                                                    const int max_input_len) {
+    // RETURN: span_mask, split_targets: for the top-down parser to fit
+    //         ldr_cache_ids : input embeddings for GPT
+    //         position ids: position id for GPT
+    //         tgt_ids for generation: cross entropy targets for GPT
+    auto buf = group_ids.request();
+    size_t buf_sz = buf.size;
+    int * group_ids_ptr = (int*)buf.ptr;
+    int group_size = group_ids_ptr[buf_sz - 1] + 1;
+
+    int max_seq_len = max_input_len * 2 - 1;
+    int max_ids_len = input_ids.shape(1);
+    auto input_ids_ptr = input_ids.unchecked<2>();
+
+    // assign best split to each cell
+    induce_best_splits(this->m_iMaxSeqLen, this->m_pCellNums, m_pCellOrders, score_splits, a_ij_splits, atom_spans);
+
+    int max_node_size = 2 * this->m_iMaxSeqLen - 1;
+
+    this->m_pLDRCache_ids = new long[group_size * max_seq_len];
+    std::fill_n(this->m_pLDRCache_ids, group_size * max_seq_len, 0);
+    this->m_pExtIds = new long[group_size * max_seq_len];
+    std::fill_n(this->m_pExtIds, group_size * max_seq_len, -1);
+    this->m_pPositionIds = new long[group_size * (max_seq_len + 1)];
+    std::fill_n(this->m_pPositionIds, group_size * (max_seq_len + 1), 0);
+    this->m_pTgtIds = new long[group_size * (max_seq_len + 1)];
+    std::fill_n(this->m_pTgtIds, group_size * (max_seq_len + 1), -1);
+
+    int mask_mat_size = (max_input_len - 1) * (max_input_len - 1);
+
+    // this->m_pSpanGatherIds = new long[group_size * max_input_len];
+    // std::fill_n(this->m_pSpanGatherIds, group_size * max_input_len, 0);
+    this->m_pSpanMasks = new long[group_size * mask_mat_size];
+    std::fill_n(this->m_pSpanMasks, group_size * mask_mat_size, 0);
+    this->m_pSplitTargets = new long[group_size * (max_input_len - 1)];
+    std::fill_n(this->m_pSplitTargets, group_size * (max_input_len - 1), -1);
+    this->m_pTokenPositions = new long[group_size * max_input_len];
+    std::fill_n(this->m_pTokenPositions, group_size * max_input_len, max_seq_len - 1);
+
+
+    Cell * queue[max_node_size];
+    int prev_group_id = -1;
+    int ldr_offset = 0;
+    int sent_offset = 1;
+    auto eos_labels_arr = eos_labels.unchecked<1>();
+    for (int batch_i = 0; batch_i < this->m_iBatchSize; ++batch_i) {
+        CellTable * tbl = this->m_pCellTables[batch_i];
+        Cell * root = tbl->get(0, tbl->getLen() - 1);
+        int sent_len = root->j - root->i + 1;
+
+        int index_bias = sent_len * 2 - 1;
+        int queue_offset = 0;
+        int split_idx = 0;
+        int group_id = group_ids_ptr[batch_i];
+        if (group_id != prev_group_id) {
+            ldr_offset = 0;
+            sent_offset = 1;
+            prev_group_id = group_id;
+        }
+
+        Cell * current = NULL;
+        queue[queue_offset++] = root;
+        while (queue_offset > 0) {
+            current = queue[--queue_offset];
+
+            if (current->split_size > 0) {
+                // non-terminal
+                int k = current->getBestSplit();
+
+                std::fill_n(m_pSpanMasks + group_id * mask_mat_size + (sent_offset - 1 + split_idx) * (max_input_len - 1) + sent_offset - 1 + current->i, current->j - current->i, 1);
+                assert(current->i <= k && current->j > k);
+                m_pSplitTargets[group_id * (max_input_len - 1) + sent_offset + split_idx - 1] = sent_offset - 1 + k;
+                ++split_idx;
+                queue[queue_offset++] = tbl->get(current->i, k);
+                queue[queue_offset++] = tbl->get(k + 1, current->j);
+            }
+        }
+
+        queue_offset = 0;
+        queue[queue_offset++] = root;
+        current = NULL;
+        // If there is appending sentence, eos_id will be overrided.
+        // tgt_ids.index({group_id, ldr_offset + index_bias}) = eos_id;
+        this->m_pTgtIds[group_id * (max_seq_len + 1) + ldr_offset + index_bias] = eos_labels_arr[group_id];
+        // for (int idx = 0; idx < sent_len - 1; ++idx) {
+        //     m_pSpanGatherIds[group_id * max_input_len + sent_offset + idx - 1] = batch_i * (max_ids_len - 1) + idx;
+        // }
+        while (queue_offset > 0) {
+            current = queue[--queue_offset];
+
+            // ldr_cache_ids.index({group_id, ldr_offset + index_bias - 1}) = current->cache_id;
+            // position_ids.index({group_id, ldr_offset + index_bias}) = current->j + sent_offset;
+            m_pLDRCache_ids[group_id * max_seq_len + ldr_offset + index_bias - 1] = current->cache_id;
+            m_pPositionIds[group_id * (max_seq_len + 1) + ldr_offset + index_bias] = current->j + sent_offset;
+            m_pExtIds[group_id * max_seq_len + ldr_offset + index_bias - 1] = current->ext_vocab_id;
+            if (current->split_size > 0) {
+                // non-terminal
+                int k = current->getBestSplit();
+
+                assert(current->i <= k && current->j > k);
+                queue[queue_offset++] = tbl->get(current->i, k);
+                queue[queue_offset++] = tbl->get(k + 1, current->j);
+                //The right node will be visited first
+
+                // input will be shift right during training, so just keep id same
+                // tgt_ids.index({group_id, ldr_offset + index_bias - 1}) = reduce_id;
+                m_pTgtIds[group_id * (max_seq_len + 1) + ldr_offset + index_bias - 1] = reduce_id;
+            } else {
+                // terminal
+                m_pTgtIds[group_id * (max_seq_len + 1) + ldr_offset + index_bias - 1] = \
+                    input_ids_ptr(batch_i, current->i);
+                m_pTokenPositions[group_id * max_input_len + sent_offset - 1 + current->i] = \
+                    ldr_offset + index_bias - 1;
+            }
+            --index_bias;
+        }
+        ldr_offset += 2 * sent_len - 1;
+        sent_offset += sent_len;
+    }
+
+    //torch::zeros({group_size, max_seq_len}, at::kLong);
+    at::Tensor ldr_cache_ids = torch::from_blob(m_pLDRCache_ids, {group_size, max_seq_len}, at::kLong);
+    // at::Tensor position_ids = torch::zeros({group_size, max_seq_len + 1}, at::kLong);
+    at::Tensor position_ids = torch::from_blob(m_pPositionIds, {group_size, max_seq_len + 1}, at::kLong);
+    // at::Tensor tgt_ids = torch::full({group_size, max_seq_len + 1}, -1, at::kLong);
+    at::Tensor tgt_ids = torch::from_blob(m_pTgtIds, {group_size, max_seq_len + 1}, at::kLong);
+    
+    at::Tensor span_masks = torch::from_blob(m_pSpanMasks, {group_size, max_input_len - 1, max_input_len - 1}, at::kLong);
+
+    at::Tensor split_targets = torch::from_blob(m_pSplitTargets, {group_size, max_input_len - 1}, at::kLong);
+
+    // at::Tensor gather_ids = torch::from_blob(m_pSpanGatherIds, {group_size, max_input_len}, at::kLong);
+    at::Tensor ext_vocab_ids_ = torch::from_blob(m_pExtIds, {group_size, max_seq_len}, at::kLong);
+
+    at::Tensor token_indices = torch::from_blob(m_pTokenPositions, {group_size, max_input_len}, at::kLong);
+
+    return {span_masks, split_targets, ldr_cache_ids, position_ids, tgt_ids, token_indices, ext_vocab_ids_};
+}
+
+WordTreeNode::WordTreeNode(int entry_id, int total_size, int depth): m_iValue(entry_id), m_iTotalSize(total_size), m_iWordId(-1), m_iDepth(depth){
+}
+
+WordTreeNode::~WordTreeNode() {
+    for (auto iter = m_mSubNodes.begin(); iter != m_mSubNodes.end(); ++iter) {
+        delete iter->second;
+    }
+}
+
+void WordTreeNode::add_ids(int * ids_ptr, int ids_len, int entry_id, int offset) {
+    int current_id = ids_ptr[offset];
+    if (m_mSubNodes.find(current_id) == m_mSubNodes.end()) {
+        m_mSubNodes[current_id] = new WordTreeNode(-1, m_iTotalSize, m_iDepth + 1);
+    }
+
+    if (offset + 1 < ids_len) {
+        m_mSubNodes[current_id]->add_ids(ids_ptr, ids_len, entry_id, offset + 1);
+    } else if(offset + 1 == ids_len) {
+        m_mSubNodes[current_id]->setWordId(entry_id);
+    }
+}
+
+WordTreeNode * WordTreeNode::next_node(int current_id) {
+    // assert(current_id >= 0 && current_id < m_iTotalSize);
+    if (m_mSubNodes.find(current_id) != m_mSubNodes.end()) {
+        return m_mSubNodes[current_id];
+    } else {
+        return NULL;
+    }
+}
+
+void WordTreeNode::setWordId(const int wordId) {
+    m_iWordId = wordId;
+}
+
+int WordTreeNode::getWordId() const {
+    return m_iWordId;
+}
+
+int WordTreeNode::getDepth() const {
+    return m_iDepth;
+}
+
+bool WordTreeNode::isWord() const {
+    return m_iWordId != -1;
+}
+
+SpanTokenizer::SpanTokenizer(vector<py::array_t<int>>& dictionary, int max_entry_id) {
+    m_pRoot = new WordTreeNode(-1, max_entry_id, false);
+    int entry_id = 0;
+    for (py::array_t<int>& ids : dictionary) {
+        auto buf = ids.request();
+        int * ids_ptr = (int*)buf.ptr;
+        int ids_len = ids.shape(0);
+        m_pRoot->add_ids(ids_ptr, ids_len, entry_id);
+        ++entry_id;
+    }
+}
+
+SpanTokenizer::~SpanTokenizer() {
+    delete m_pRoot;
+}
+
+vector<int> SpanTokenizer::tokenize(py::array_t<int>& ids_arr) {
+    auto buf = ids_arr.request();
+    int * ids_ptr = (int*)buf.ptr;
+    int size = ids_arr.shape(0);
+
+    list<WordTreeNode*> * current_nodes = new list<WordTreeNode*>();
+    list<WordTreeNode*> * next_nodes = new list<WordTreeNode*>();
+    list<WordTreeNode*> * temp_nodes = 0;
+    current_nodes->push_back(m_pRoot);
+    WordTreeNode * temp = NULL;
+    vector<int> results;
+    for (int idx = 0; idx < size; ++idx) {
+        next_nodes->push_back(m_pRoot);
+        for (WordTreeNode * activated_node : *current_nodes) {
+            temp = activated_node->next_node(ids_ptr[idx]);
+            if (temp != NULL) {
+                next_nodes->push_back(temp);
+                if (temp->isWord()) {
+                    // TODO: record span position and word id
+                    results.push_back(idx);
+                    results.push_back(temp->getDepth());
+                    results.push_back(temp->getWordId());
+                    // temp->getWordId();
+                }
+            }
+        }
+
+        temp_nodes = current_nodes;
+        current_nodes = next_nodes;
+        next_nodes = temp_nodes;
+        next_nodes->clear();
+    }
+    delete current_nodes;
+    delete next_nodes;
+    return results;
+}

gpst/cpp_extension/py_backend.h

@@ -0,0 +1,204 @@
+// Copyright (c) 2024 Ant Group
+// Author: Xiang Hu
+#pragma once
+#include <vector>
+#include <torch/extension.h>
+#include <pybind11/pybind11.h>
+#include <pybind11/numpy.h>
+#include <pybind11/stl.h>
+#include <list>
+#include <map>
+using namespace std;
+namespace py = pybind11;
+
+
+class Cell;
+class TableManager;
+
+struct Span {
+    int i;
+    int j;
+};
+
+class LinkedNode {
+private:
+    Cell * m_pCell;
+    LinkedNode * m_pLeft;
+    LinkedNode * m_pRight;
+    LinkedNode * m_pLeftup;
+    LinkedNode * m_pRightup;
+    LinkedNode * m_pLeftdown;
+    LinkedNode * m_pRightdown;
+public:
+    LinkedNode(Cell * value);
+    ~LinkedNode();
+
+    Cell * getCell() const;
+    LinkedNode * left() const;
+    LinkedNode * right() const;
+    LinkedNode * leftup() const;
+    LinkedNode * rightup() const;
+    LinkedNode * leftdown() const;
+    LinkedNode * rightdown() const;
+
+    void setLeft(LinkedNode * other);
+    void setRight(LinkedNode * other);
+    void setLeftup(LinkedNode * other);
+    void setRightup(LinkedNode * other);
+    void setLeftdown(LinkedNode * other);
+    void setRightdown(LinkedNode * other);
+};
+
+
+class Cell {
+private:
+    LinkedNode * m_pNode;
+    list<Cell*> m_lParents;
+    TableManager * m_pMgr;
+    bool m_bIsRoot;
+    int m_pReadyChild;
+
+    void notifyChildReady(); // notify when one of its inside cell is ready
+public:
+    const int i;
+    const int j;
+    const int batch_id;
+    int ext_vocab_id;
+    int cache_id;
+    int * splits;
+    int split_size;
+    int best_split;
+    int a_ij_split;
+    bool detached;
+    
+    Cell(int i, int j, int window_size, TableManager * mgr, const int batch_id, const bool is_root);
+    ~Cell();
+
+    int getDetachedCacheID(int detach_offset) const;
+    LinkedNode * getNode() const;
+    void setNode(LinkedNode * target);
+    int getBestSplit() const ;
+    int getGumbelSplit() const ;
+    void addParent(Cell * parent);
+    void onReady(); // call when a cell is ready to encode
+};
+
+class CellTable {
+private:
+    Cell ** m_pCells;
+    Cell ** m_pCreatedCells;
+    TableManager * m_pMgr;
+    const int m_iSeqLen;
+    const int m_iBatchId;
+    const int m_iMaxCreatedCells;
+    const int m_iWindowSize;
+    int m_iCellOffset;
+public:
+    CellTable(int seq_len, int window_size, int batch_i, TableManager * mgr);
+    ~CellTable();
+
+    int getLen() const;
+
+    Cell * get(const int i, const int j); // get cell at i,j, if null then create one
+    bool isEmpty(const int i, const int j);
+};
+
+// class ActiveCells {
+// public:
+//     ActiveCells(int window_size, int seq_len, CellTable * cell_table);
+//     ~ActiveCells(); //delete memory
+
+//     vector<LinkedNode*> prune(LinkedNode * node); //return created Nodes
+// }
+
+// init_active_cells(int window_size, int seq_len, CellTable * cell_table);
+// vector<LinkedNode*> prune(LinkedNode * node); //return created Nodes
+
+class TableManager {
+private:
+    CellTable ** m_pCellTables;
+    // Span ** m_pMergeOrders;
+    int m_iCurrentStep;
+    int m_iCellNum;
+    const int m_iCacheOffset;
+    const int m_iDetachCacheOffset;
+    const int m_iWindowSize;
+    const int m_iBatchSize;
+    int m_iMaxSeqLen;
+
+    Cell *** m_pCellOrders;
+    int * m_pCellNums;
+    long ** m_pTargetCacheIds;
+    long ** m_pGroupCacheIds;
+    long ** m_pTargetExtIds;
+    long ** m_pDetachGroupCacheIds;
+
+    long * m_pLDRCache_ids;
+    long * m_pPositionIds;
+    long * m_pExtIds; // external vocab id for each position
+    long * m_pTgtIds;
+
+    long * m_pSpanMasks;
+    long * m_pSplitTargets;
+    long * m_pSpanGatherIds;
+    long * m_pTokenPositions;
+    
+    list<Cell*> m_lReadyCells;
+private:
+    void push_cell(Cell * cell);
+    void build_cell_dependencies(Span ** pMergeOrders);
+public:
+    TableManager(const py::array_t<int>& seq_lens, const py::array_t<int>& group_ids,
+                 const py::array_t<int>& merge_orders, const int window_size, 
+                 const int cache_id_offset, const int detach_cache_id_offset,
+                 vector<py::array_t<int>>& span_ids);
+    ~TableManager();
+    bool is_finished();
+    vector<at::Tensor> step();
+    // vector<at::Tensor> best_trees(py::array_t<int>& best_splits);
+
+    vector<at::Tensor> prepare_generation(vector<py::array_t<int>>& score_splits, 
+                                          vector<py::array_t<int>>& a_ij_splits, 
+                                          vector<py::array_t<int>>& atom_spans, 
+                                          const py::array_t<int>& input_ids,
+                                          const py::array_t<int>& groups_ids, 
+                                          const py::array_t<int>& eos_labels,
+                                          const int reduce_id,
+                                          const int max_input_len);
+    at::Tensor root_ids();
+    at::Tensor prepare_bilm(int total_len, int bos_id, int eos_id);
+    const int batch_size() const;
+    void on_cell_ready(Cell* cell);
+};
+
+class WordTreeNode;
+
+class WordTreeNode {
+private:
+    // WordTreeNode ** m_pSubNodes;
+    map<int, WordTreeNode*> m_mSubNodes;
+    const int m_iTotalSize;
+    const int m_iValue;
+    int m_iWordId;
+    const int m_iDepth;
+public:
+    WordTreeNode(int entry_id, int total_size, int depth=0);
+    void add_ids(int * ids_ptr, int ids_len, int entry_id, int offset=0);
+    WordTreeNode * next_node(int current_id);
+    ~WordTreeNode();
+    void setWordId(const int wordId);
+    int getWordId() const;
+    int getDepth() const;
+    bool isWord() const;
+    void print_path() const;
+};
+
+class SpanTokenizer {
+private:
+    WordTreeNode * m_pRoot;
+public:
+    SpanTokenizer(vector<py::array_t<int>>& dictionary, int max_entry_id);
+    ~SpanTokenizer();
+
+    vector<int> tokenize(py::array_t<int>& ids_arr);
+};

gpst/data_structure/py_backend.py

@@ -0,0 +1,74 @@
+# coding=utf-8
+# Copyright (c) 2024 Ant Group
+# Author: Xiang Hu
+from typing import List
+import torch
+import numpy as np
+from gpst import cppbackend
+    
+class CPPChartTableManager:
+    def __init__(self, seq_lens, window_size, merge_orders, cache_id_offset, detach_cache_id_offset, group_ids=None, span_ids=None):
+        # seq_lens: np array
+        # merge_orders: np array
+        self.seq_lens = seq_lens
+        
+        if group_ids is None:
+            group_ids = [i for i in range(len(seq_lens))]
+        if span_ids is None:
+            span_ids = []
+        assert len(group_ids) == len(seq_lens)
+        self.cpp_tbl_mgr = cppbackend.TableManager(seq_lens, group_ids, merge_orders, window_size, 
+                                                   cache_id_offset, detach_cache_id_offset, span_ids)
+        self._root_ids = None
+    
+    @property
+    def root_ids(self):
+        return self._root_ids
+    
+    def construct_inside_groups(self, device):
+        target_cache_ids_list = []
+        span_ids_batch_list = []
+        cache_groups_batch_list = []
+        detach_cache_groups_batch_list = []
+        total_time = None
+        while not self.cpp_tbl_mgr.is_finished():
+            tgt_cache_ids, span_ids, cache_ids, detach_cache_ids = self.cpp_tbl_mgr.step()
+            target_cache_ids_list.append(tgt_cache_ids.to(device, non_blocking=True))
+            span_ids_batch_list.append(span_ids.to(device, non_blocking=True))
+            cache_groups_batch_list.append(cache_ids.to(device, non_blocking=True))
+            detach_cache_groups_batch_list.append(detach_cache_ids.to(device, non_blocking=True))
+        
+        self._root_ids = self.cpp_tbl_mgr.root_ids().to(device, non_blocking=True)
+        return target_cache_ids_list, span_ids_batch_list, cache_groups_batch_list, detach_cache_groups_batch_list
+    
+    # def best_trees(self, best_splits, atom_spans=None, terminal_only=False):
+    #     if atom_spans is None:
+    #         atom_spans = [torch.zeros((0,2))] * self.cpp_tbl_mgr.batch_size()
+    #     else:
+    #         atom_spans = [torch.tensor(spans) if len(spans) > 0 else torch.zeros((0, 2)) for spans in atom_spans]
+    #     assert len(atom_spans) == self.cpp_tbl_mgr.batch_size()
+    #     # np_arr = [t.data.numpy() for t in best_splits]
+    #     # best_splits = np.concatenate(np_arr)
+    #     # return targets, cache_ids
+    #     splits, cache_ids = self.cpp_tbl_mgr.best_trees(best_splits, atom_spans, terminal_only)
+    #     return splits, cache_ids
+
+    def prepare_generation(self, score_orders, split_orders, atom_spans, input_ids, groups_ids, eos_id, reduce_id, max_input_len,
+                           eos_labels=None):
+        if atom_spans is None:
+            atom_spans = [np.zeros((0,2))] * self.cpp_tbl_mgr.batch_size()
+        else:
+            atom_spans = [np.array(spans) if len(spans) > 0 else np.zeros((0, 2)) for spans in atom_spans]
+        assert len(atom_spans) == self.cpp_tbl_mgr.batch_size()
+        score_orders = [order.data.numpy() for order in score_orders]
+        split_orders = [order.data.numpy() for order in split_orders]
+        if eos_labels is None:
+            eos_labels = np.full((groups_ids[-1] + 1), fill_value=eos_id)
+        assert len(eos_labels) == groups_ids[-1] + 1
+        span_masks, split_targets, ldr_cache_ids, position_ids, tgt_ids, token_indices, ext_ids = \
+            self.cpp_tbl_mgr.prepare_generation(score_orders, split_orders, atom_spans, input_ids, groups_ids, 
+                                                eos_labels, reduce_id, max_input_len)
+        return span_masks, split_targets, ldr_cache_ids, position_ids, tgt_ids, token_indices, ext_ids
+    
+    def prepare_bilm(self, total_len, bos_id, eos_id):
+        return self.cpp_tbl_mgr.prepare_bilm(total_len, bos_id, eos_id)

gpst/data_structure/tensor_cache.py

@@ -0,0 +1,121 @@
+# coding=utf-8
+# Copyright (c) 2024 Ant Group
+# Author: Xiang Hu
+from enum import Enum
+import torch
+
+
+class CacheType(Enum):
+    NORMAL = 0
+    DETACH = 1
+
+
+class TensorCache:
+    def __init__(self, max_window,
+                 seq_lens,
+                 cache_types,
+                 dims,
+                 placeholder_num,
+                 device, 
+                 iter_times=1,
+                 total_cache_size=-1):
+        self.placeholder_num = placeholder_num  #
+        self.total_block_size = placeholder_num
+        self.iter_times = iter_times
+
+        if total_cache_size == -1:
+            self.block_sizes = [0] * len(seq_lens)
+            self._init_blocks_size(seq_lens, max_window)
+            for i in range(len(seq_lens)):
+                self.total_block_size += self.block_sizes[i]
+        else:
+            self.total_block_size += total_cache_size
+
+        self._max_lengths = [placeholder_num] * len(cache_types)
+        self._cache_num = len(cache_types)
+        self.cache_types = cache_types
+        self.caches = [None] * self._cache_num
+        self.dims = dims
+        self.device = device
+        # dtype = torch.float16 if torch.is_autocast_enabled() else torch.float
+        for i, cache_type in enumerate(cache_types):
+            if cache_type == CacheType.NORMAL:
+                self.caches[i] = torch.full((self.total_block_size, dims[i]), 0.0, dtype=torch.float32, device=device)
+            elif cache_type == CacheType.DETACH:
+                self.caches[i] = torch.full((self.total_block_size * 2, dims[i]), 0.0, dtype=torch.float32, device=device)
+
+    @property
+    def capacity(self):
+        return self.total_block_size
+
+    @property
+    def detach_offset(self):
+        return self.total_block_size
+
+    def init_placeholders(self, cache_ids, values):
+        for cache_id, value in zip(cache_ids, values):
+            self.caches[cache_id][:self.placeholder_num] = value
+
+    def _init_blocks_size(self, seq_lens, max_window):
+        seq_num = len(seq_lens)
+        for seq_i in range(seq_num):
+            seq_len = seq_lens[seq_i]
+            block_max_len = 0
+            for layer_i in range(seq_len):
+                if layer_i <= max_window:
+                    block_max_len += (seq_len - layer_i) * (layer_i + 1)
+                else:
+                    block_max_len += (max_window + 1) * max_window
+            self.block_sizes[seq_i] = block_max_len * self.iter_times
+
+    def gather(self, indices, cache_ids):
+        # Gather tensors according to CacheItem pairs
+        tensors_gathered = []
+        if isinstance(indices, torch.Tensor):
+            gather_indices = indices
+        else:
+            gather_indices = torch.tensor(indices, dtype=torch.long, device=self.device)
+        for cache_id in cache_ids:
+            tensor_block = self.caches[cache_id]
+            tensor_gather = tensor_block.index_select(dim=0, index=gather_indices)
+            tensors_gathered.append(tensor_gather)
+        return tensors_gathered
+
+    def fill(self, cache_id_offset, cache_id_len, cache_ids, values):
+        if len(cache_ids) != len(values):
+            raise Exception('TensorCache::fill names and values mismatch')
+        # cdef PyObject ** tensor_block
+        for cache_id, value in zip(cache_ids, values):
+            tensor_block = self.caches[cache_id]
+            tensor_block[cache_id_offset: cache_id_offset + cache_id_len] = value
+            if self.cache_types[cache_id] == CacheType.DETACH:
+                detach_offset = self.total_block_size + cache_id_offset
+                tensor_block[detach_offset: detach_offset + cache_id_len] = value.detach()
+
+    def get(self, cache_id, idx):
+        assert self.caches[cache_id] is not None
+        return self.caches[cache_id][idx]
+
+    def get_tensor_cache(self, cache_id):
+        return self.caches[cache_id]
+
+    def detach(self, idx):
+        detach_idx = self.total_block_size + idx
+        for cache_i, cache_type in enumerate(self.cache_types):
+            if cache_type == CacheType.DETACH:
+                self.caches[cache_i][detach_idx] = self.caches[cache_i][idx].detach()
+        return detach_idx
+
+    def scatter(self, indices, cache_ids, values):
+        if isinstance(indices, torch.Tensor):
+            scatter_indices = indices
+        else:
+            scatter_indices = torch.tensor(indices, dtype=torch.long, device=self.device)
+        
+        for cache_id, value in zip(cache_ids, values):
+            tensor_block = self.caches[cache_id]
+            dim = value.shape[-1]
+            scatter_indices_ = scatter_indices.unsqueeze(1).repeat(1, dim)
+            if value.dtype != tensor_block.dtype:
+                value = value.to(tensor_block.dtype)
+            tensor_block.scatter_(dim=0, index=scatter_indices_, src=value)

gpst/fast_parser.py

@@ -0,0 +1,103 @@
+# coding=utf-8
+# Copyright (c) 2024 Ant Group
+# Author: Xiang Hu
+from gpst.topdown_parser import BasicParser
+from gpst.Llama_flash_attn import *
+import torch.nn as nn
+import torch
+import numpy as np
+
+
+class TransformerParser(BasicParser):
+    def __init__(self, config) -> None:
+        super().__init__()
+        self.legacy_mode = False
+        self.hidden_dim = config.parser_hidden_dim
+        self.input_dim = config.parser_input_dim
+
+        self.score_mlp = nn.Sequential(nn.Linear(2 * self.input_dim, self.hidden_dim),
+                                       nn.GELU(),
+                                       nn.Dropout(config.hidden_dropout_prob),
+                                       nn.Linear(self.hidden_dim, 1))
+        # self.score_mlp = nn.Sequential(nn.Linear(self.input_dim, self.hidden_dim),
+        #                                nn.GELU(),
+        #                                nn.Dropout(config.hidden_dropout_prob),
+        #                                nn.Linear(self.hidden_dim, 1))
+
+        # args = ModelArgs(config.parser_input_dim, config.parser_num_layers, config.parser_nhead,
+        #                  config.vocab_size, max_seq_len=config.parser_max_len, apply_norm=False)
+        args = ModelArgs(config.parser_input_dim, config.parser_num_layers, config.parser_nhead,
+                         config.vocab_size, max_seq_len=config.parser_max_len, apply_norm=True)
+
+
+        # layer = nn.TransformerEncoderLayer(self.input_dim, nhead=config.parser_nhead,
+        #                                    dim_feedforward=self.hidden_dim, activation='gelu',
+        #                                    batch_first=True)
+        # self.encoder = nn.TransformerEncoder(layer, config.parser_num_layers)
+        self.encoder = Transformer(args)
+
+    def _generate_flatten_input_ids(self, input_ids, attn_mask, group_ids):
+        seq_lens = attn_mask.sum(dim=1).cpu().data.numpy()
+        batch_size = group_ids[-1] + 1
+        group_lengths = [0] * batch_size
+        for sent_id, group_id in enumerate(group_ids):
+            group_lengths[group_id] += seq_lens[sent_id]
+
+        max_length = max(group_lengths)
+
+        prev_group_id = -1
+        flatten_ids = input_ids.new_zeros((batch_size, max_length))
+        flatten_masks = attn_mask.new_zeros([batch_size, max_length])
+        for sent_id, group_id in enumerate(group_ids):
+            if prev_group_id != group_id:
+                offset = 0
+                prev_group_id = group_id
+            flatten_ids[group_id, offset: offset + seq_lens[sent_id]] = input_ids[sent_id, :seq_lens[sent_id]]
+            flatten_masks[group_id, offset: offset + seq_lens[sent_id]] = 1
+            offset += seq_lens[sent_id]
+        return flatten_ids, flatten_masks, seq_lens
+
+    def _recover_score_chunks(self, org_shape, scores, seq_lens, group_ids):
+        rev_scores = scores.new_zeros((org_shape[0], org_shape[1] - 1))  # (N, L)
+        offset = 0
+        prev_group_id = -1
+        for sent_id, group_id in enumerate(group_ids):
+            if group_id != prev_group_id:
+                prev_group_id = group_id
+                offset = 0
+            sent_len = seq_lens[sent_id]
+            rev_scores[sent_id, : sent_len - 1] = scores[group_id, offset: offset + sent_len - 1]
+            offset += sent_len
+        return rev_scores
+
+    def _split_point_scores(self, input_ids, attn_mask, group_ids=None):
+        # attn_mask: (N, L) recording segment ids
+        # if group_ids is not None:
+        #     # reorgniaze input_ids
+        #     org_input_ids, org_mask = input_ids, attn_mask
+        #     input_ids, attn_mask, seq_lens = self._generate_flatten_input_ids(input_ids, attn_mask, group_ids)
+        if attn_mask is None:
+            attn_mask = torch.ones_like(input_ids)
+        # print(attn_mask.shape)
+        attn_mask = attn_mask.unsqueeze(2) == attn_mask.unsqueeze(1)  # (N, L, L) or (L, L)
+        # print(attn_mask.shape)
+        mask = torch.zeros_like(attn_mask, dtype=torch.float)
+        mask.masked_fill_(attn_mask == 0, -np.inf)
+        # if len(attn_mask.shape) == 3:
+        #     eye_mask = torch.eye(attn_mask.shape[1], device=input_ids.device)
+        #     mask.masked_fill_(attn_mask + eye_mask.unsqueeze(0) == 0, -np.inf)
+        # else:
+        #     mask.masked_fill_(attn_mask == 0, -np.inf)
+        # seq_lens = attn_mask.sum(dim=-1)  # (N)
+        N = input_ids.shape[0]
+        pos_ids = torch.arange(input_ids.shape[1], device=input_ids.device)
+        outputs = self.encoder(input_ids, attn_mask=mask, position_ids=pos_ids)
+        split_logits = torch.cat([outputs[:, :-1, :], outputs[:, 1:, :]], dim=-1)  # (N, L - 1, 2 * dim)
+        # dim = outputs.shape[-1]
+        # split_logits = torch.cat([outputs[:, :-1, dim//2:], outputs[:, 1:, :dim//2]], dim=-1)
+        scores = self.score_mlp(split_logits)
+        scores = scores.squeeze(-1)
+        # if group_ids is not None:
+        #     # split scores
+        #     scores = self._recover_score_chunks(org_input_ids.shape, scores, seq_lens, group_ids)
+        return scores

gpst/generative_r2d2_fast.py

@@ -0,0 +1,299 @@
+# coding=utf-8
+# Copyright (c) 2024 Ant Group
+# Author: Xiang Hu
+import torch.nn as nn
+import torch
+import torch.nn.functional as F
+from datetime import datetime
+from concurrent.futures import ThreadPoolExecutor
+from dataclasses import dataclass
+from typing import Optional
+from transformers import PreTrainedModel, PretrainedConfig, GPT2Config
+from transformers.modeling_outputs import MaskedLMOutput
+from gpst.gpt2_flash_attn import GPT2Model
+from gpst.r2d2_insideoutside import *
+import copy
+
+def load_model(model, model_path, strict=True):
+    state_dict = torch.load(model_path, map_location=lambda a, b: a)
+    transfered_state_dict = {}
+    for k, v in state_dict.items():
+        new_k = k.replace('module.', '')
+        transfered_state_dict[new_k] = v
+    model.load_state_dict(transfered_state_dict, strict=strict)
+
+def index_sanity_hook(module, input, output):
+    def check(tensor):
+        if isinstance(tensor, torch.Tensor):
+            if tensor.dtype == torch.long or "int" in str(tensor.dtype):
+                if tensor.max() > 10000 or tensor.min() < -10000:  # change threshold if needed
+                    print(f"[!] Suspicious index in {module.__class__.__name__}: min={tensor.min().item()}, max={tensor.max().item()}, shape={tensor.shape}")
+
+    # Check all inputs
+    for item in input:
+        if isinstance(item, (tuple, list)):
+            for sub in item:
+                check(sub)
+        else:
+            check(item)
+
+@dataclass(kw_only=True)
+class R2D2GenOutput():
+    struct_loss: Optional[torch.FloatTensor] = None,
+    non_struct_loss: Optional[torch.FloatTensor] = None,
+    non_struct_loss_fullscale: Optional[torch.FloatTensor] = None,
+    action_logits: Optional[torch.FloatTensor] = None,
+    hidden_states: Optional[torch.FloatTensor] = None,
+    cls_hidden_states: Optional[torch.FloatTensor] = None,
+    tgt_ids: Optional[torch.LongTensor] = None, 
+    pred: Optional[torch.FloatTensor] = None, 
+    splits: Optional[torch.LongTensor] = None,
+    gpt_loss: Optional[torch.FloatTensor] = None,
+    action_loss: Optional[torch.FloatTensor] = None,
+    inside_outside_loss: Optional[torch.FloatTensor] = None,
+    parser_loss: Optional[torch.FloatTensor] = None, 
+    glue_finetune_loss: Optional[torch.FloatTensor] = None,
+    past_kv: Optional[torch.FloatTensor] = None
+    logits: Optional[torch.FloatTensor] = None,
+    loss: Optional[torch.FloatTensor] = None,
+
+
+class GPSTConfig(PretrainedConfig):
+    model_type = "gpst"
+    
+    def __init__(self, r2d2=None, gpt=None, **kwargs):#, gptconfig, r2d2config, **kwargs):
+
+        self.gptconfig = gpt
+        self.r2d2config = r2d2
+        super().__init__(**kwargs)
+
+class GPST(PreTrainedModel):
+    config_class = GPSTConfig
+
+    def __init__(self, config, gradient_checkpoint=False):
+        super().__init__(config)
+        self.config = config
+        self.r2d2_config = PretrainedConfig.from_dict(config.r2d2config)
+        self.gpt_config = GPT2Config.from_dict(config.gptconfig)
+
+        self.vocab_size = self.gpt_config.vocab_size
+
+        total_layer = self.gpt_config.n_layer
+        action_transformers = GPT2Model(copy.deepcopy(self.gpt_config), no_embedding=True, no_layer_norm=True, n_layers_manual=self.gpt_config.action_layer_num)
+        action_transformers.gradient_checkpointing = gradient_checkpoint
+        self.gpt_config.n_layer = total_layer - self.gpt_config.action_layer_num
+        self.gpt_config.num_hidden_layers = total_layer - self.gpt_config.action_layer_num
+        gpt_transformers = GPT2Model(self.gpt_config, no_embedding=True, no_extra_embedding=True)
+        gpt_transformers.gradient_checkpointing = gradient_checkpoint
+
+        r2d2 = InsideOutsideModule(self.r2d2_config)
+        self.model = FastGenerativeR2D2(
+            r2d2=r2d2, 
+            action_layers=action_transformers, 
+            generation_layers=gpt_transformers, 
+            vocab_size=self.vocab_size,
+            r2d2_input_dim=r2d2.input_dim,
+            embedding_dim=self.gpt_config.n_embd,
+            ext_vocab_size=self.r2d2_config.ext_vocab_size,
+            dense_hidden_factor=self.gpt_config.dense_hidden_factor
+        )
+
+        #for name, module in self.model.named_modules():
+        #    #module.register_forward_hook(index_sanity_hook)
+        #    module.register_full_backward_hook(index_sanity_hook)
+
+    def get_input_embeddings(self):
+        return self.model.embeddings
+    
+    def forward(self, **kwargs):
+        return self.model(**kwargs)
+
+class FastGenerativeR2D2(nn.Module):
+    def __init__(self, r2d2, action_layers, generation_layers, vocab_size, 
+                 r2d2_input_dim, embedding_dim, dropout_rate=0.2, ext_vocab_size=0, 
+                 fix_embeddings=False, dense_hidden_factor=4):
+        # embedding dim is used to feed to r2d2
+        # input dim is sued to feed to GPT
+        super().__init__()
+        self.embedding_dim = embedding_dim  # embedding_dim > r2d2_input_dim
+        self.r2d2_input_dim = r2d2_input_dim
+        self.r2d2 = r2d2
+
+        self.vocab_size = vocab_size
+
+        # self.action_ln = nn.Linear(self.embedding_dim, 2)  # judge reduce or predict next token
+
+        self.enable_gpt = False
+        if action_layers is not None and generation_layers is not None:
+            self.dense_hidden_factor = dense_hidden_factor
+            self.action_layers = action_layers
+            self.generation_layers = generation_layers
+            self.bos_embedding = nn.Parameter(torch.rand(self.embedding_dim))
+            self.up_scale = nn.Linear(self.r2d2_input_dim, self.embedding_dim)
+            self.dense = nn.Sequential(nn.Linear(self.embedding_dim, self.dense_hidden_factor * self.embedding_dim),
+                                        nn.GELU(),
+                                        nn.Dropout(dropout_rate),
+                                        nn.Linear(self.dense_hidden_factor * self.embedding_dim, self.embedding_dim))
+            self.action_mlp = nn.Sequential(nn.LayerNorm(self.embedding_dim),
+                                nn.Linear(self.embedding_dim, self.embedding_dim),
+                                nn.GELU(),
+                                nn.Dropout(dropout_rate),
+                                nn.Linear(self.embedding_dim, 2))
+            self.enable_gpt = True
+        
+        self.classifier = nn.Linear(self.embedding_dim, vocab_size, bias=False)
+        self.embeddings = nn.Embedding(vocab_size, self.embedding_dim)
+        self.embeddings.requires_grad = not fix_embeddings
+        self.down_scale = nn.Linear(self.embedding_dim, self.r2d2_input_dim)
+
+        self.insideoutside_dense = nn.Sequential(
+            nn.Linear(r2d2_input_dim, self.dense_hidden_factor * r2d2_input_dim),
+            nn.GELU(),
+            nn.Dropout(dropout_rate),
+            nn.Linear(self.dense_hidden_factor * r2d2_input_dim, self.embedding_dim)
+        )
+
+        # self.parallel_stream = torch.cuda.Stream()
+
+        self._init_weights()
+        self._tie_weights()
+
+    def _init_weights(self):
+        if self.enable_gpt:
+            self.bos_embedding.data.normal_(mean=0, std=0.02)
+        self.embeddings.weight.data.normal_(mean=0, std=0.02)
+
+    def _tie_weights(self):
+        self.classifier.weight = self.embeddings.weight
+
+    def get_parser(self):
+        return self.r2d2.parser
+        
+    def from_pretrain(self, model_path, strict=True):
+        load_model(self, model_path, strict=strict)
+        self._tie_weights()
+
+    def _append_eos_label(self, eos_labels, chunk_input_ids, chunk_masks, next_token_indices, max_input_len):
+        chunk_masks = (chunk_masks.sum(dim=1) > 0).to(int)
+        seq_lens = chunk_masks.sum(dim=1)  # (N)
+        temp_ids = torch.zeros((chunk_input_ids.shape[0], chunk_input_ids.shape[1] + 1), dtype=chunk_input_ids.dtype, device=chunk_input_ids.device)
+        temp_ids.fill_(-100)
+        temp_ids[:, :-1] = chunk_input_ids
+        # comment this line to support discriminant way
+        temp_ids.scatter_(1, seq_lens.unsqueeze(1), torch.tensor(eos_labels, device=chunk_input_ids.device).unsqueeze(1))
+        chunk_input_ids = temp_ids
+        next_token_indices = next_token_indices[:, :max_input_len + 1]
+        return next_token_indices, chunk_input_ids
+
+    def forward(self, chunk_input_ids= None, chunk_masks=None, input_ids=None, masks=None, eos_labels=None, group_ids=None, 
+                atom_spans=None, span_ids=None, external_vocab_ids=None, 
+                coeff=1.0, temperature=1.0, past_key_values=None):
+        batch_size = max(group_ids) + 1
+        r2d2_input_ids = torch.where(chunk_input_ids == -100, 0, chunk_input_ids)
+        input_embeddings = self.embeddings(r2d2_input_ids)
+        r2d2_embeddings = self.down_scale(input_embeddings)
+        # max_input_len = chunk_input_ids.shape[1]
+        max_input_len = (chunk_masks != 0).sum(dim=1).max().to('cpu', non_blocking=True)
+        
+        ctx, outside_tgt, ldr_repr, position_ids, tgt_ids, token_indices, ext_ids, split_targets, l_height = \
+            self.r2d2(r2d2_input_ids, chunk_masks, input_ids, masks, r2d2_embeddings, group_ids, 
+                      max_input_len, atom_spans=atom_spans, coeff=coeff, temperature=temperature, span_ids=span_ids,
+                      eos_labels=eos_labels, external_vocab_ids=external_vocab_ids)
+
+
+        if self.training:
+            # with torch.cuda.stream(self.parallel_stream):
+            parser_loss = self.r2d2.parser_loss(ctx)
+            outside_embeddings = self.r2d2.outside_embeddings(ctx) # (num non-padding tokens in batch) x embedding_dim
+            io_dense = self.insideoutside_dense(outside_embeddings) #  (num non-padding tokens in batch) x embedding_dim
+            outside_logits = self.classifier(io_dense) # (num non-padding tokens in batch) x voc_size
+            insideoutside_loss = F.cross_entropy(outside_logits, outside_tgt)
+        else:
+            parser_loss = insideoutside_loss = 0
+        
+        logits = action_logits = None
+        gpt_loss = action_loss = 0
+        past_kv = None
+        hidden_states = None
+
+        if self.enable_gpt:
+            if past_key_values is not None:
+                action_past_kv, gen_past_kv = past_key_values
+            else:
+                action_past_kv = gen_past_kv = None
+            gpt_input = self.up_scale(ldr_repr).clone() # ldr_repr: batch_size x (2*max_seq_len - 1) x HP dim; gpt_input: batch_size x  (2*max_seq_len - 1) x emb_dim 
+            gpt_input.scatter_(1, token_indices.unsqueeze(2).repeat(1, 1, input_embeddings.shape[-1]).clone(), 
+                                input_embeddings.to(gpt_input.dtype)) # inserting values of input_embeddings at token_indices
+            
+            # ext_embedding = self.ext_embeds(ext_ids)
+            # gpt_input = gpt_input + ext_embedding
+            bos_emb = self.bos_embedding.unsqueeze(0).repeat(batch_size, 1)
+            # position ids already considered <bos>
+            cat_input = torch.cat([bos_emb.unsqueeze(1), gpt_input], dim=1)  # batch_size x 2*max_seq_len x emb_dim # old comment:  (group_size, L + 1, dim) where L is (2*max_seq_len - 1)
+            # cat_input = self.layer_norm(cat_input)
+            # cat_input = self.norm(cat_input)
+            outputs = self.action_layers(inputs_embeds=cat_input, position_ids=position_ids, past_key_values=action_past_kv)  # (B, L, dim)
+            action_logits = self.action_mlp(outputs.last_hidden_state)  # (batch_size x 2*max_seq_len x emb_dim) to (B, 2*max_seq_len, 2)
+            # before, tgt_ids has shape (batch_size x 2*max_seq_len) and has the reduce_token_id (default: 50257) in some placdse. Paddingn: -1
+            # first where expression: action_tgt has shape (batch_size x 2*max_seq_len), 0 and 1 values for generate/reduce. Padding is still zero
+            # second where: padding with -1 
+            action_tgt = torch.where(tgt_ids == self.r2d2.reduce_id, 1, 0)  # REDUCE: 1, SHIFT:0
+            action_tgt = torch.where(tgt_ids != -1, action_tgt, -1)
+            # print(action_tgt)
+
+            next_token_indices = (tgt_ids != self.r2d2.reduce_id).int().argsort(dim=-1, descending=True, stable=True)  # (B, L)
+            if eos_labels is None:
+                #truncated_len = max_input_len if self.training else max_input_len + 1
+                truncated_len = max_input_len
+                next_token_indices = next_token_indices[:, :truncated_len] # batch_size x longest_sequence;  
+            else:
+                next_token_indices, chunk_input_ids = self._append_eos_label(eos_labels, chunk_input_ids, chunk_masks, next_token_indices, max_input_len)
+            # outputs.last_hidden_state: batch_size x 2*max_seq_len x HP dim
+            # generation_inputs: batch_size x max_seq_len x HP dim 
+            # next_token_indices_reformat: batch_size x max_seq_len x HP dim
+            next_token_indices_reformat = next_token_indices.unsqueeze(2).repeat(1, 1, self.embedding_dim)
+            generation_inputs = outputs.last_hidden_state.gather(1, next_token_indices_reformat) 
+            # token_pos_ids = position_ids.gather(1, next_token_indices)
+            # gather outputs to predict the next token
+            # token_outputs: LM output. Last hidden_state has batch_size x max_seq_len x HP dim
+            token_outputs = self.generation_layers(inputs_embeds=generation_inputs, past_key_values=gen_past_kv)
+
+            hidden_states = token_outputs.last_hidden_state
+            logits = self.classifier(self.dense(hidden_states))  # new: batch_size x max_seq_len x voc_size old: (group_size, L + 1, vocab)
+            # predict token loss + action loss
+            # print("chunk_input_ids: ", chunk_input_ids)
+            #if self.training:
+            gpt_loss = F.cross_entropy(logits.permute(0, 2, 1), chunk_input_ids, ignore_index=-100)
+            action_loss = F.cross_entropy(action_logits.permute(0, 2, 1), action_tgt, ignore_index=-1)
+            past_kv = (outputs.past_key_values, token_outputs.past_key_values)
+
+        # torch.cuda.synchronize()
+        # return loss + lm_loss + parser_loss, split_targets
+        return R2D2GenOutput(struct_loss=insideoutside_loss + l_height, 
+                             non_struct_loss=0.5 * gpt_loss + action_loss + parser_loss,
+                             non_struct_loss_fullscale=gpt_loss + action_loss + parser_loss,
+                             logits=logits,
+                             action_logits=action_logits,
+                             hidden_states=hidden_states, 
+                             tgt_ids=chunk_input_ids, 
+                             gpt_loss=gpt_loss,
+                             action_loss=action_loss,
+                             inside_outside_loss=insideoutside_loss,
+                             parser_loss=parser_loss,
+                             past_kv=past_kv,
+                             splits=split_targets,
+                             loss=action_loss+gpt_loss+parser_loss+insideoutside_loss+l_height)
+        # parser_loss should be fine
+
+class FastGenerativeR2D2_discriminant_glue(FastGenerativeR2D2):
+    
+    def _append_eos_label(self, eos_labels, chunk_input_ids, chunk_masks, next_token_indices, max_input_len):
+        chunk_masks = (chunk_masks.sum(dim=1) > 0).to(int)
+        seq_lens = chunk_masks.sum(dim=1)  # (N)
+        temp_ids = torch.zeros((chunk_input_ids.shape[0], chunk_input_ids.shape[1] + 1), dtype=chunk_input_ids.dtype, device=chunk_input_ids.device)
+        temp_ids.fill_(-100)
+        temp_ids[:, :-1] = chunk_input_ids
+        # temp_ids.scatter_(1, seq_lens.unsqueeze(1), torch.tensor(eos_labels, device=chunk_input_ids.device).unsqueeze(1))
+        chunk_input_ids = temp_ids
+        next_token_indices = next_token_indices[:, :max_input_len + 1]
+        return next_token_indices, chunk_input_ids

gpst/gpt2_flash_attn.py

@@ -0,0 +1,1625 @@
+"""PyTorch OpenAI GPT-2 model."""
+
+import math
+import os
+import warnings
+from dataclasses import dataclass
+from typing import Optional, Tuple, Union
+
+import torch
+import torch.utils.checkpoint
+from torch import nn
+from torch.cuda.amp import autocast
+from torch.nn import BCEWithLogitsLoss, CrossEntropyLoss, MSELoss
+import torch.nn.functional as F
+
+from transformers.activations import ACT2FN
+from transformers.modeling_outputs import (
+    BaseModelOutputWithPastAndCrossAttentions,
+    CausalLMOutputWithCrossAttentions,
+    QuestionAnsweringModelOutput,
+    SequenceClassifierOutputWithPast,
+    TokenClassifierOutput,
+)
+from transformers.modeling_utils import PreTrainedModel, SequenceSummary
+from transformers.pytorch_utils import Conv1D, find_pruneable_heads_and_indices, prune_conv1d_layer
+from transformers.utils import (
+    ModelOutput,
+    add_code_sample_docstrings,
+    add_start_docstrings,
+    add_start_docstrings_to_model_forward,
+    logging,
+    replace_return_docstrings,
+)
+from transformers.utils.model_parallel_utils import assert_device_map, get_device_map
+from transformers import GPT2Config
+
+
+logger = logging.get_logger(__name__)
+
+_CHECKPOINT_FOR_DOC = "gpt2"
+_CONFIG_FOR_DOC = "GPT2Config"
+
+GPT2_PRETRAINED_MODEL_ARCHIVE_LIST = [
+    "gpt2",
+    "gpt2-medium",
+    "gpt2-large",
+    "gpt2-xl",
+    "distilgpt2",
+    # See all GPT-2 models at https://huggingface.co/models?filter=gpt2
+]
+
+
+def load_tf_weights_in_gpt2(model, config, gpt2_checkpoint_path):
+    """Load tf checkpoints in a pytorch model"""
+    try:
+        import re
+
+        import tensorflow as tf
+    except ImportError:
+        logger.error(
+            "Loading a TensorFlow model in PyTorch, requires TensorFlow to be installed. Please see "
+            "https://www.tensorflow.org/install/ for installation instructions."
+        )
+        raise
+    tf_path = os.path.abspath(gpt2_checkpoint_path)
+    logger.info(f"Converting TensorFlow checkpoint from {tf_path}")
+    # Load weights from TF model
+    init_vars = tf.train.list_variables(tf_path)
+    names = []
+    arrays = []
+    for name, shape in init_vars:
+        logger.info(f"Loading TF weight {name} with shape {shape}")
+        array = tf.train.load_variable(tf_path, name)
+        names.append(name)
+        arrays.append(array.squeeze())
+
+    for name, array in zip(names, arrays):
+        name = name[6:]  # skip "model/"
+        name = name.split("/")
+        pointer = model
+        for m_name in name:
+            if re.fullmatch(r"[A-Za-z]+\d+", m_name):
+                scope_names = re.split(r"(\d+)", m_name)
+            else:
+                scope_names = [m_name]
+            if scope_names[0] == "w" or scope_names[0] == "g":
+                pointer = getattr(pointer, "weight")
+            elif scope_names[0] == "b":
+                pointer = getattr(pointer, "bias")
+            elif scope_names[0] == "wpe" or scope_names[0] == "wte":
+                pointer = getattr(pointer, scope_names[0])
+                pointer = getattr(pointer, "weight")
+            else:
+                pointer = getattr(pointer, scope_names[0])
+            if len(scope_names) >= 2:
+                num = int(scope_names[1])
+                pointer = pointer[num]
+        try:
+            if pointer.shape != array.shape:
+                raise ValueError(f"Pointer shape {pointer.shape} and array shape {array.shape} mismatched")
+        except ValueError as e:
+            e.args += (pointer.shape, array.shape)
+            raise
+        logger.info(f"Initialize PyTorch weight {name}")
+        pointer.data = torch.from_numpy(array)
+    return model
+
+
+class GPT2Attention(nn.Module):
+    def __init__(self, config, is_cross_attention=False, layer_idx=None):
+        super().__init__()
+
+        max_positions = config.max_position_embeddings
+        self.register_buffer(
+            "bias",
+            torch.tril(torch.ones((2 * max_positions, 2 * max_positions), dtype=torch.bool)).view(
+                1, 1, 2 * max_positions, 2 * max_positions
+            ),
+            persistent=False,
+        )
+        self.register_buffer("masked_bias", torch.tensor(-1e4), persistent=False)
+
+        self.embed_dim = config.hidden_size
+        self.num_heads = config.num_attention_heads
+        self.head_dim = self.embed_dim // self.num_heads
+        self.split_size = self.embed_dim
+        if self.head_dim * self.num_heads != self.embed_dim:
+            raise ValueError(
+                f"`embed_dim` must be divisible by num_heads (got `embed_dim`: {self.embed_dim} and `num_heads`:"
+                f" {self.num_heads})."
+            )
+
+        self.scale_attn_weights = config.scale_attn_weights
+        self.is_cross_attention = is_cross_attention
+
+        # Layer-wise attention scaling, reordering, and upcasting
+        self.scale_attn_by_inverse_layer_idx = config.scale_attn_by_inverse_layer_idx
+        self.layer_idx = layer_idx
+        self.reorder_and_upcast_attn = config.reorder_and_upcast_attn
+
+        if self.is_cross_attention:
+            self.c_attn = Conv1D(2 * self.embed_dim, self.embed_dim)
+            self.q_attn = Conv1D(self.embed_dim, self.embed_dim)
+        else:
+            self.c_attn = Conv1D(3 * self.embed_dim, self.embed_dim)
+        self.c_proj = Conv1D(self.embed_dim, self.embed_dim)
+
+        # self.attn_dropout = nn.Dropout(config.attn_pdrop)
+        self.attn_dropout = config.attn_pdrop
+        self.resid_dropout = nn.Dropout(config.resid_pdrop)
+
+        self.pruned_heads = set()
+
+    def prune_heads(self, heads):
+        if len(heads) == 0:
+            return
+        heads, index = find_pruneable_heads_and_indices(heads, self.num_heads, self.head_dim, self.pruned_heads)
+        index_attn = torch.cat([index, index + self.split_size, index + (2 * self.split_size)])
+
+        # Prune conv1d layers
+        self.c_attn = prune_conv1d_layer(self.c_attn, index_attn, dim=1)
+        self.c_proj = prune_conv1d_layer(self.c_proj, index, dim=0)
+
+        # Update hyper params
+        self.split_size = (self.split_size // self.num_heads) * (self.num_heads - len(heads))
+        self.num_heads = self.num_heads - len(heads)
+        self.pruned_heads = self.pruned_heads.union(heads)
+
+    def _attn(self, query, key, value, attention_mask=None, head_mask=None):
+        # print(f'query input: {query[:, 0, :, :3]}')
+        # L, S = query.size(-2), key.size(-2)
+        # temp_mask = torch.ones(L, S, dtype=torch.bool).tril(diagonal=0)
+        # attn_bias = torch.zeros(L, S, dtype=query.dtype)
+        # attn_bias.masked_fill_(temp_mask.logical_not(), float("-inf"))
+        # print(attn_bias)
+        with torch.backends.cuda.sdp_kernel():
+            dropout_p = self.attn_dropout if self.training else 0
+            L, S = query.size(-2), key.size(-2)
+            if attention_mask is None:
+                if L == S:
+                    attn_outputs = F.scaled_dot_product_attention(query, key, value, dropout_p=dropout_p, is_causal=True)
+                else:
+                    attn_mask = torch.ones(L, S, dtype=torch.bool, device=query.device).tril(diagonal=S - L)
+                    # # print(attn_mask)
+                    attn_outputs = F.scaled_dot_product_attention(query, key, value, dropout_p=dropout_p, attn_mask=attn_mask)
+            else:
+                attn_outputs = F.scaled_dot_product_attention(query, key, value, dropout_p=dropout_p, attn_mask=attention_mask)
+            # elif query.size(-2) == 1:
+            #     # inference mode
+            #     attn_outputs = F.scaled_dot_product_attention(query, key, value, dropout_p=dropout_p, attn_mask=attention_mask)
+            # else:
+            #     raise Exception('unsupported mode')
+        # print(attn_outputs[:, :, :3])
+        # print(attn_outputs[:, 0, :, :3])
+        return attn_outputs
+        # attn_weights = torch.matmul(query, key.transpose(-1, -2))
+
+        # if self.scale_attn_weights:
+        #     attn_weights = attn_weights / torch.full(
+        #         [], value.size(-1) ** 0.5, dtype=attn_weights.dtype, device=attn_weights.device
+        #     )
+
+        # # Layer-wise attention scaling
+        # if self.scale_attn_by_inverse_layer_idx:
+        #     attn_weights = attn_weights / float(self.layer_idx + 1)
+
+        # if not self.is_cross_attention:
+        #     # if only "normal" attention layer implements causal mask
+        #     query_length, key_length = query.size(-2), key.size(-2)
+        #     causal_mask = self.bias[:, :, key_length - query_length : key_length, :key_length]
+        #     mask_value = torch.finfo(attn_weights.dtype).min
+        #     # Need to be a tensor, otherwise we get error: `RuntimeError: expected scalar type float but found double`.
+        #     # Need to be on the same device, otherwise `RuntimeError: ..., x and y to be on the same device`
+        #     mask_value = torch.full([], mask_value, dtype=attn_weights.dtype).to(attn_weights.device)
+        #     attn_weights = torch.where(causal_mask, attn_weights.to(attn_weights.dtype), mask_value)
+
+        # if attention_mask is not None:
+        #     # Apply the attention mask
+        #     attn_weights = attn_weights + attention_mask
+
+        # attn_weights = nn.functional.softmax(attn_weights, dim=-1)
+
+        # # Downcast (if necessary) back to V's dtype (if in mixed-precision) -- No-Op otherwise
+        # attn_weights = attn_weights.type(value.dtype)
+        # attn_weights = self.attn_dropout(attn_weights)
+
+        # # Mask heads if we want to
+        # if head_mask is not None:
+        #     attn_weights = attn_weights * head_mask
+
+        # attn_output = torch.matmul(attn_weights, value)
+
+        # return attn_output, attn_weights
+
+    def _upcast_and_reordered_attn(self, query, key, value, attention_mask=None, head_mask=None):
+        # Use `torch.baddbmm` (a bit more efficient w/ alpha param for scaling -- from Megatron-LM)
+        raise Exception('not implemented for GPT2WithFlashAttn')
+
+    def _split_heads(self, tensor, num_heads, attn_head_size):
+        """
+        Splits hidden_size dim into attn_head_size and num_heads
+        """
+        new_shape = tensor.size()[:-1] + (num_heads, attn_head_size)
+        tensor = tensor.view(new_shape)
+        return tensor.permute(0, 2, 1, 3)  # (batch, head, seq_length, head_features)
+
+    def _merge_heads(self, tensor, num_heads, attn_head_size):
+        """
+        Merges attn_head_size dim and num_attn_heads dim into hidden_size
+        """
+        tensor = tensor.permute(0, 2, 1, 3).contiguous()
+        new_shape = tensor.size()[:-2] + (num_heads * attn_head_size,)
+        return tensor.view(new_shape)
+
+    def forward(
+        self,
+        hidden_states: Optional[Tuple[torch.FloatTensor]],
+        layer_past: Optional[Tuple[torch.Tensor]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        encoder_hidden_states: Optional[torch.Tensor] = None,
+        encoder_attention_mask: Optional[torch.FloatTensor] = None,
+        use_cache: Optional[bool] = False,
+        output_attentions: Optional[bool] = False,
+    ) -> Tuple[Union[torch.Tensor, Tuple[torch.Tensor]], ...]:
+        if encoder_hidden_states is not None:
+            if not hasattr(self, "q_attn"):
+                raise ValueError(
+                    "If class is used as cross attention, the weights `q_attn` have to be defined. "
+                    "Please make sure to instantiate class with `GPT2Attention(..., is_cross_attention=True)`."
+                )
+
+            query = self.q_attn(hidden_states)
+            key, value = self.c_attn(encoder_hidden_states).split(self.split_size, dim=2)
+            attention_mask = encoder_attention_mask
+        else:
+            query, key, value = self.c_attn(hidden_states).split(self.split_size, dim=2)
+
+        query = self._split_heads(query, self.num_heads, self.head_dim)
+        key = self._split_heads(key, self.num_heads, self.head_dim)
+        value = self._split_heads(value, self.num_heads, self.head_dim)
+
+        if layer_past is not None:
+            past_key, past_value = layer_past
+            key = torch.cat((past_key, key), dim=-2)
+            value = torch.cat((past_value, value), dim=-2)
+
+        if use_cache is True:
+            present = (key, value)
+        else:
+            present = None
+
+        attn_output = self._attn(query, key, value, attention_mask, head_mask)
+
+        attn_output = self._merge_heads(attn_output, self.num_heads, self.head_dim)
+        attn_output = self.c_proj(attn_output)
+        attn_output = self.resid_dropout(attn_output)
+
+        outputs = (attn_output, present)
+
+        return outputs  # a, present, (attentions)
+
+
+class GPT2MLP(nn.Module):
+    def __init__(self, intermediate_size, config):
+        super().__init__()
+        embed_dim = config.hidden_size
+        self.c_fc = Conv1D(intermediate_size, embed_dim)
+        self.c_proj = Conv1D(embed_dim, intermediate_size)
+        self.act = ACT2FN[config.activation_function]
+        self.dropout = nn.Dropout(config.resid_pdrop)
+
+    def forward(self, hidden_states: Optional[Tuple[torch.FloatTensor]]) -> torch.FloatTensor:
+        hidden_states = self.c_fc(hidden_states)
+        hidden_states = self.act(hidden_states)
+        hidden_states = self.c_proj(hidden_states)
+        hidden_states = self.dropout(hidden_states)
+        return hidden_states
+
+
+class GPT2Block(nn.Module):
+    def __init__(self, config, layer_idx=None):
+        super().__init__()
+        hidden_size = config.hidden_size
+        inner_dim = config.n_inner if config.n_inner is not None else 4 * hidden_size
+
+        self.ln_1 = nn.LayerNorm(hidden_size, eps=config.layer_norm_epsilon)
+        self.attn = GPT2Attention(config, layer_idx=layer_idx)
+        self.ln_2 = nn.LayerNorm(hidden_size, eps=config.layer_norm_epsilon)
+
+        if config.add_cross_attention:
+            self.crossattention = GPT2Attention(config, is_cross_attention=True, layer_idx=layer_idx)
+            self.ln_cross_attn = nn.LayerNorm(hidden_size, eps=config.layer_norm_epsilon)
+
+        self.mlp = GPT2MLP(inner_dim, config)
+
+    def forward(
+        self,
+        hidden_states: Optional[Tuple[torch.FloatTensor]],
+        layer_past: Optional[Tuple[torch.Tensor]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        encoder_hidden_states: Optional[torch.Tensor] = None,
+        encoder_attention_mask: Optional[torch.FloatTensor] = None,
+        use_cache: Optional[bool] = False,
+        output_attentions: Optional[bool] = False,
+    ) -> Union[Tuple[torch.Tensor], Optional[Tuple[torch.Tensor, Tuple[torch.FloatTensor, ...]]]]:
+        residual = hidden_states
+        hidden_states = self.ln_1(hidden_states)
+        attn_outputs = self.attn(
+            hidden_states,
+            layer_past=layer_past,
+            attention_mask=attention_mask,
+            head_mask=head_mask,
+            use_cache=use_cache,
+            output_attentions=output_attentions,
+        )
+        attn_output = attn_outputs[0]  # output_attn: a, present, (attentions)
+        outputs = attn_outputs[1:]
+        # residual connection
+        hidden_states = attn_output + residual
+
+        if encoder_hidden_states is not None:
+            # add one self-attention block for cross-attention
+            if not hasattr(self, "crossattention"):
+                raise ValueError(
+                    f"If `encoder_hidden_states` are passed, {self} has to be instantiated with "
+                    "cross-attention layers by setting `config.add_cross_attention=True`"
+                )
+            residual = hidden_states
+            hidden_states = self.ln_cross_attn(hidden_states)
+            cross_attn_outputs = self.crossattention(
+                hidden_states,
+                attention_mask=attention_mask,
+                head_mask=head_mask,
+                encoder_hidden_states=encoder_hidden_states,
+                encoder_attention_mask=encoder_attention_mask,
+                output_attentions=output_attentions,
+            )
+            attn_output = cross_attn_outputs[0]
+            # residual connection
+            hidden_states = residual + attn_output
+            outputs = outputs + cross_attn_outputs[2:]  # add cross attentions if we output attention weights
+
+        residual = hidden_states
+        hidden_states = self.ln_2(hidden_states)
+        feed_forward_hidden_states = self.mlp(hidden_states)
+        # residual connection
+        hidden_states = residual + feed_forward_hidden_states
+
+        if use_cache:
+            outputs = (hidden_states,) + outputs
+        else:
+            outputs = (hidden_states,) + outputs[1:]
+
+        return outputs  # hidden_states, present, (attentions, cross_attentions)
+
+
+class GPT2PreTrainedModel(PreTrainedModel):
+    """
+    An abstract class to handle weights initialization and a simple interface for downloading and loading pretrained
+    models.
+    """
+
+    config_class = GPT2Config
+    load_tf_weights = load_tf_weights_in_gpt2
+    base_model_prefix = "transformer"
+    is_parallelizable = True
+    supports_gradient_checkpointing = True
+    _no_split_modules = ["GPT2Block"]
+    _skip_keys_device_placement = "past_key_values"
+
+    def __init__(self, *inputs, **kwargs):
+        super().__init__(*inputs, **kwargs)
+
+    def _init_weights(self, module):
+        """Initialize the weights."""
+        if isinstance(module, (nn.Linear, Conv1D)):
+            # Slightly different from the TF version which uses truncated_normal for initialization
+            # cf https://github.com/pytorch/pytorch/pull/5617
+            module.weight.data.normal_(mean=0.0, std=self.config.initializer_range)
+            if module.bias is not None:
+                module.bias.data.zero_()
+        elif isinstance(module, nn.Embedding):
+            module.weight.data.normal_(mean=0.0, std=self.config.initializer_range)
+            if module.padding_idx is not None:
+                module.weight.data[module.padding_idx].zero_()
+        elif isinstance(module, nn.LayerNorm):
+            module.bias.data.zero_()
+            module.weight.data.fill_(1.0)
+
+        # Reinitialize selected weights subject to the OpenAI GPT-2 Paper Scheme:
+        #   > A modified initialization which accounts for the accumulation on the residual path with model depth. Scale
+        #   > the weights of residual layers at initialization by a factor of 1/√N where N is the # of residual layers.
+        #   >   -- GPT-2 :: https://openai.com/blog/better-language-models/
+        #
+        # Reference (Megatron-LM): https://github.com/NVIDIA/Megatron-LM/blob/main/megatron/model/gpt_model.py
+        for name, p in module.named_parameters():
+            if name == "c_proj.weight":
+                # Special Scaled Initialization --> There are 2 Layer Norms per Transformer Block
+                p.data.normal_(mean=0.0, std=(self.config.initializer_range / math.sqrt(2 * self.config.n_layer)))
+
+    def _set_gradient_checkpointing(self, module, value=False):
+        if isinstance(module, GPT2Model):
+            module.gradient_checkpointing = value
+
+
+@dataclass
+class GPT2DoubleHeadsModelOutput(ModelOutput):
+    """
+    Base class for outputs of models predicting if two sentences are consecutive or not.
+
+    Args:
+        loss (`torch.FloatTensor` of shape `(1,)`, *optional*, returned when `labels` is provided):
+            Language modeling loss.
+        mc_loss (`torch.FloatTensor` of shape `(1,)`, *optional*, returned when `mc_labels` is provided):
+            Multiple choice classification loss.
+        logits (`torch.FloatTensor` of shape `(batch_size, num_choices, sequence_length, config.vocab_size)`):
+            Prediction scores of the language modeling head (scores for each vocabulary token before SoftMax).
+        mc_logits (`torch.FloatTensor` of shape `(batch_size, num_choices)`):
+            Prediction scores of the multiple choice classification head (scores for each choice before SoftMax).
+        past_key_values (`Tuple[Tuple[torch.Tensor]]`, *optional*, returned when `use_cache=True` is passed or when `config.use_cache=True`):
+            Tuple of length `config.n_layers`, containing tuples of tensors of shape `(batch_size, num_heads,
+            sequence_length, embed_size_per_head)`).
+
+            Contains pre-computed hidden-states (key and values in the attention blocks) that can be used (see
+            `past_key_values` input) to speed up sequential decoding.
+        hidden_states (`tuple(torch.FloatTensor)`, *optional*, returned when `output_hidden_states=True` is passed or when `config.output_hidden_states=True`):
+            Tuple of `torch.FloatTensor` (one for the output of the embeddings + one for the output of each layer) of
+            shape `(batch_size, sequence_length, hidden_size)`.
+
+            Hidden-states of the model at the output of each layer plus the initial embedding outputs.
+        attentions (`tuple(torch.FloatTensor)`, *optional*, returned when `output_attentions=True` is passed or when `config.output_attentions=True`):
+            Tuple of `torch.FloatTensor` (one for each layer) of shape `(batch_size, num_heads, sequence_length,
+            sequence_length)`.
+
+            GPT2Attentions weights after the attention softmax, used to compute the weighted average in the
+            self-attention heads.
+    """
+
+    loss: Optional[torch.FloatTensor] = None
+    mc_loss: Optional[torch.FloatTensor] = None
+    logits: torch.FloatTensor = None
+    mc_logits: torch.FloatTensor = None
+    past_key_values: Optional[Tuple[Tuple[torch.FloatTensor]]] = None
+    hidden_states: Optional[Tuple[torch.FloatTensor]] = None
+    attentions: Optional[Tuple[torch.FloatTensor]] = None
+
+
+GPT2_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 ([`GPT2Config`]): 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.
+"""
+
+GPT2_INPUTS_DOCSTRING = r"""
+    Args:
+        input_ids (`torch.LongTensor` of shape `(batch_size, input_ids_length)`):
+            `input_ids_length` = `sequence_length` if `past_key_values` is `None` else
+            `past_key_values[0][0].shape[-2]` (`sequence_length` of input past key value states). Indices of input
+            sequence tokens in the vocabulary.
+
+            If `past_key_values` is used, only `input_ids` that do not have their past calculated should be passed as
+            `input_ids`.
+
+            Indices can be obtained using [`AutoTokenizer`]. See [`PreTrainedTokenizer.encode`] and
+            [`PreTrainedTokenizer.__call__`] for details.
+
+            [What are input IDs?](../glossary#input-ids)
+        past_key_values (`Tuple[Tuple[torch.Tensor]]` of length `config.n_layers`):
+            Contains precomputed hidden-states (key and values in the attention blocks) as computed by the model (see
+            `past_key_values` output below). Can be used to speed up sequential decoding. The `input_ids` which have
+            their past given to this model should not be passed as `input_ids` as they have already been computed.
+        attention_mask (`torch.FloatTensor` 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**.
+
+            If `past_key_values` is used, `attention_mask` needs to contain the masking strategy that was used for
+            `past_key_values`. In other words, the `attention_mask` always has to have the length:
+            `len(past_key_values) + len(input_ids)`
+
+            [What are attention masks?](../glossary#attention-mask)
+        token_type_ids (`torch.LongTensor` of shape `(batch_size, input_ids_length)`, *optional*):
+            Segment token indices to indicate first and second portions of the inputs. Indices are selected in `[0,
+            1]`:
+
+            - 0 corresponds to a *sentence A* token,
+            - 1 corresponds to a *sentence B* token.
+
+            [What are token type IDs?](../glossary#token-type-ids)
+        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.max_position_embeddings - 1]`.
+
+            [What are position IDs?](../glossary#position-ids)
+        head_mask (`torch.FloatTensor` of shape `(num_heads,)` or `(num_layers, num_heads)`, *optional*):
+            Mask to nullify selected heads of the self-attention modules. Mask values selected in `[0, 1]`:
+
+            - 1 indicates the head is **not masked**,
+            - 0 indicates the head is **masked**.
+
+        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.
+
+            If `past_key_values` is used, optionally only the last `inputs_embeds` have to be input (see
+            `past_key_values`).
+        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.
+"""
+PARALLELIZE_DOCSTRING = r"""
+    This is an experimental feature and is a subject to change at a moment's notice.
+
+    Uses a device map to distribute attention modules of the model across several devices. If no device map is given,
+    it will evenly distribute blocks across all devices.
+
+    Args:
+        device_map (`Dict[int, list]`, optional, defaults to None):
+            A dictionary that maps attention modules to devices. Note that the embedding module and LMHead are always
+            automatically mapped to the first device (for esoteric reasons). That means that the first device should
+            have fewer attention modules mapped to it than other devices. For reference, the gpt2 models have the
+            following number of attention modules:
+
+                - gpt2: 12
+                - gpt2-medium: 24
+                - gpt2-large: 36
+                - gpt2-xl: 48
+
+    Example:
+
+    ```python
+    # Here is an example of a device map on a machine with 4 GPUs using gpt2-xl, which has a total of 48 attention modules:
+    model = GPT2LMHeadModel.from_pretrained("gpt2-xl")
+    device_map = {
+        0: [0, 1, 2, 3, 4, 5, 6, 7, 8],
+        1: [9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21],
+        2: [22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34],
+        3: [35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47],
+    }
+    model.parallelize(device_map)
+    ```
+"""
+DEPARALLELIZE_DOCSTRING = r"""
+    Moves the model to cpu from a model parallel state.
+
+    Example:
+
+    ```python
+    # On a 4 GPU machine with gpt2-large:
+    model = GPT2LMHeadModel.from_pretrained("gpt2-large")
+    device_map = {
+        0: [0, 1, 2, 3, 4, 5, 6, 7],
+        1: [8, 9, 10, 11, 12, 13, 14, 15],
+        2: [16, 17, 18, 19, 20, 21, 22, 23],
+        3: [24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35],
+    }
+    model.parallelize(device_map)  # Splits the model across several devices
+    model.deparallelize()  # Put the model back on cpu and cleans memory by calling torch.cuda.empty_cache()
+    ```
+"""
+
+
+@add_start_docstrings(
+    "The bare GPT2 Model transformer outputting raw hidden-states without any specific head on top.",
+    GPT2_START_DOCSTRING,
+)
+class GPT2Model(GPT2PreTrainedModel):
+    def __init__(self, config, no_embedding=False, no_extra_embedding=False, no_layer_norm=False, n_layers_manual=None):
+        super().__init__(config)
+
+        if n_layers_manual:
+            self.config.num_hidden_layers = n_layers_manual
+            self.config.n_layer = n_layers_manual
+
+        self.embed_dim = config.hidden_size
+
+        if not no_embedding:
+            self.wte = nn.Embedding(config.vocab_size, self.embed_dim)
+
+        if not no_extra_embedding:
+            self.wpe = nn.Embedding(config.max_position_embeddings, self.embed_dim)
+
+        self.drop = nn.Dropout(config.embd_pdrop)
+        self.h = nn.ModuleList([GPT2Block(config, layer_idx=i) for i in range(config.num_hidden_layers)])
+
+        if not no_layer_norm:
+            self.ln_f = nn.LayerNorm(self.embed_dim, eps=config.layer_norm_epsilon)
+        self.no_extra_embedding = no_extra_embedding
+        self.no_layer_norm = no_layer_norm
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+        self.gradient_checkpointing = False
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+
+    def get_input_embeddings(self):
+        return self.wte
+
+    def set_input_embeddings(self, new_embeddings):
+        self.wte = new_embeddings
+
+    def _prune_heads(self, heads_to_prune):
+        """
+        Prunes heads of the model. heads_to_prune: dict of {layer_num: list of heads to prune in this layer}
+        """
+        for layer, heads in heads_to_prune.items():
+            self.h[layer].attn.prune_heads(heads)
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING)
+    @add_code_sample_docstrings(
+        checkpoint=_CHECKPOINT_FOR_DOC,
+        output_type=BaseModelOutputWithPastAndCrossAttentions,
+        config_class=_CONFIG_FOR_DOC,
+    )
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Tuple[Tuple[torch.Tensor]]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        encoder_hidden_states: Optional[torch.Tensor] = None,
+        encoder_attention_mask: 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, BaseModelOutputWithPastAndCrossAttentions]:
+        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
+
+        if input_ids is not None and inputs_embeds is not None:
+            raise ValueError("You cannot specify both input_ids and inputs_embeds at the same time")
+        elif input_ids is not None:
+            input_shape = input_ids.size()
+            input_ids = input_ids.view(-1, input_shape[-1])
+            batch_size = input_ids.shape[0]
+        elif inputs_embeds is not None:
+            input_shape = inputs_embeds.size()[:-1]
+            batch_size = inputs_embeds.shape[0]
+        else:
+            raise ValueError("You have to specify either input_ids or inputs_embeds")
+
+        device = input_ids.device if input_ids is not None else inputs_embeds.device
+
+        if token_type_ids is not None:
+            token_type_ids = token_type_ids.view(-1, input_shape[-1])
+        if position_ids is not None:
+            position_ids = position_ids.view(-1, input_shape[-1])
+
+        if past_key_values is None:
+            past_length = 0
+            past_key_values = tuple([None] * len(self.h))
+        else:
+            past_length = past_key_values[0][0].size(-2)
+        if position_ids is None:
+            position_ids = torch.arange(past_length, input_shape[-1] + past_length, dtype=torch.long, device=device)
+            position_ids = position_ids.unsqueeze(0).view(-1, input_shape[-1])
+
+        # GPT2Attention mask.
+        if attention_mask is not None:
+            if batch_size <= 0:
+                raise ValueError("batch_size has to be defined and > 0")
+            attention_mask = attention_mask.view(batch_size, -1)
+            # We create a 3D attention mask from a 2D tensor mask.
+            # Sizes are [batch_size, 1, 1, to_seq_length]
+            # So we can broadcast to [batch_size, num_heads, from_seq_length, to_seq_length]
+            # this attention mask is more simple than the triangular masking of causal attention
+            # used in OpenAI GPT, we just need to prepare the broadcast dimension here.
+            attention_mask = attention_mask[:, None, None, :]
+
+            # Since attention_mask is 1.0 for positions we want to attend and 0.0 for
+            # masked positions, this operation will create a tensor which is 0.0 for
+            # positions we want to attend and the dtype's smallest value for masked positions.
+            # Since we are adding it to the raw scores before the softmax, this is
+            # effectively the same as removing these entirely.
+            attention_mask = attention_mask.to(dtype=self.dtype)  # fp16 compatibility
+            attention_mask = (1.0 - attention_mask) * torch.finfo(self.dtype).min
+
+        # If a 2D or 3D attention mask is provided for the cross-attention
+        # we need to make broadcastable to [batch_size, num_heads, seq_length, seq_length]
+        if self.config.add_cross_attention and encoder_hidden_states is not None:
+            encoder_batch_size, encoder_sequence_length, _ = encoder_hidden_states.size()
+            encoder_hidden_shape = (encoder_batch_size, encoder_sequence_length)
+            if encoder_attention_mask is None:
+                encoder_attention_mask = torch.ones(encoder_hidden_shape, device=device)
+            encoder_attention_mask = self.invert_attention_mask(encoder_attention_mask)
+        else:
+            encoder_attention_mask = None
+
+        # Prepare head mask if needed
+        # 1.0 in head_mask indicate we keep the head
+        # attention_probs has shape bsz x n_heads x N x N
+        # head_mask has shape n_layer x batch x n_heads x N x N
+        head_mask = self.get_head_mask(head_mask, self.config.n_layer)
+
+        if inputs_embeds is None:
+            inputs_embeds = self.wte(input_ids)
+        
+        if not self.no_extra_embedding:
+            position_embeds = self.wpe(position_ids)
+            hidden_states = inputs_embeds + position_embeds
+
+            if token_type_ids is not None:
+                token_type_embeds = self.wte(token_type_ids)
+                hidden_states = hidden_states + token_type_embeds
+        else:
+            hidden_states = inputs_embeds
+
+        hidden_states = self.drop(hidden_states)
+
+        output_shape = (-1,) + input_shape[1:] + (hidden_states.size(-1),)
+
+        if self.gradient_checkpointing and self.training:
+            if use_cache:
+                # logger.warning(
+                #     "`use_cache=True` is incompatible with gradient checkpointing. Setting `use_cache=False`..."
+                # )
+                pass
+
+        presents = () if use_cache else None
+        all_self_attentions = () if output_attentions else None
+        all_cross_attentions = () if output_attentions and self.config.add_cross_attention else None
+        all_hidden_states = () if output_hidden_states else None
+        for i, (block, layer_past) in enumerate(zip(self.h, past_key_values)):
+            # Model parallel
+            if self.model_parallel:
+                torch.cuda.set_device(hidden_states.device)
+                # Ensure layer_past is on same device as hidden_states (might not be correct)
+                if layer_past is not None:
+                    layer_past = tuple(past_state.to(hidden_states.device) for past_state in layer_past)
+                # Ensure that attention_mask is always on the same device as hidden_states
+                if attention_mask is not None:
+                    attention_mask = attention_mask.to(hidden_states.device)
+                if isinstance(head_mask, torch.Tensor):
+                    head_mask = head_mask.to(hidden_states.device)
+            if output_hidden_states:
+                all_hidden_states = all_hidden_states + (hidden_states,)
+
+            if self.gradient_checkpointing and self.training:
+
+                def create_custom_forward(module):
+                    def custom_forward(*inputs):
+                        # None for past_key_value
+                        return module(*inputs, use_cache, output_attentions)
+
+                    return custom_forward
+
+                outputs = torch.utils.checkpoint.checkpoint(
+                    create_custom_forward(block),
+                    hidden_states,
+                    None,
+                    attention_mask,
+                    head_mask[i],
+                    encoder_hidden_states,
+                    encoder_attention_mask,
+                )
+            else:
+                outputs = block(
+                    hidden_states,
+                    layer_past=layer_past,
+                    attention_mask=attention_mask,
+                    head_mask=head_mask[i],
+                    encoder_hidden_states=encoder_hidden_states,
+                    encoder_attention_mask=encoder_attention_mask,
+                    use_cache=use_cache,
+                    output_attentions=output_attentions,
+                )
+
+            hidden_states = outputs[0]
+            if use_cache is True:
+                presents = presents + (outputs[1],)
+
+            if output_attentions:
+                all_self_attentions = all_self_attentions + (outputs[2 if use_cache else 1],)
+                if self.config.add_cross_attention:
+                    all_cross_attentions = all_cross_attentions + (outputs[3 if use_cache else 2],)
+
+            # Model Parallel: If it's the last layer for that device, put things on the next device
+            if self.model_parallel:
+                for k, v in self.device_map.items():
+                    if i == v[-1] and "cuda:" + str(k) != self.last_device:
+                        hidden_states = hidden_states.to("cuda:" + str(k + 1))
+
+        if not self.no_layer_norm:
+            hidden_states = self.ln_f(hidden_states)
+
+        hidden_states = hidden_states.view(output_shape)
+        # Add last hidden state
+        if output_hidden_states:
+            all_hidden_states = all_hidden_states + (hidden_states,)
+
+        if not return_dict:
+            return tuple(
+                v
+                for v in [hidden_states, presents, all_hidden_states, all_self_attentions, all_cross_attentions]
+                if v is not None
+            )
+
+        return BaseModelOutputWithPastAndCrossAttentions(
+            last_hidden_state=hidden_states,
+            past_key_values=presents,
+            hidden_states=all_hidden_states,
+            attentions=all_self_attentions,
+            cross_attentions=all_cross_attentions,
+        )
+
+
+@add_start_docstrings(
+    """
+    The GPT2 Model transformer with a language modeling head on top (linear layer with weights tied to the input
+    embeddings).
+    """,
+    GPT2_START_DOCSTRING,
+)
+class GPT2LMHeadModel(GPT2PreTrainedModel):
+    _tied_weights_keys = ["lm_head.weight"]
+
+    def __init__(self, config):
+        super().__init__(config)
+        self.transformer = GPT2Model(config)
+        self.lm_head = nn.Linear(config.n_embd, config.vocab_size, bias=False)
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    @add_start_docstrings(PARALLELIZE_DOCSTRING)
+    def parallelize(self, device_map=None):
+        warnings.warn(
+            "`GPT2LMHeadModel.parallelize` is deprecated and will be removed in v5 of Transformers, you should load"
+            " your model with `device_map='balanced'` in the call to `from_pretrained`. You can also provide your own"
+            " `device_map` but it needs to be a dictionary module_name to device, so for instance {'transformer.h.0':"
+            " 0, 'transformer.h.1': 1, ...}",
+            FutureWarning,
+        )
+        self.device_map = (
+            get_device_map(len(self.transformer.h), range(torch.cuda.device_count()))
+            if device_map is None
+            else device_map
+        )
+        assert_device_map(self.device_map, len(self.transformer.h))
+        self.transformer.parallelize(self.device_map)
+        self.lm_head = self.lm_head.to(self.transformer.first_device)
+        self.model_parallel = True
+
+    @add_start_docstrings(DEPARALLELIZE_DOCSTRING)
+    def deparallelize(self):
+        warnings.warn(
+            "Like `parallelize`, `deparallelize` is deprecated and will be removed in v5 of Transformers.",
+            FutureWarning,
+        )
+        self.transformer.deparallelize()
+        self.transformer = self.transformer.to("cpu")
+        self.lm_head = self.lm_head.to("cpu")
+        self.model_parallel = False
+        torch.cuda.empty_cache()
+
+    def get_output_embeddings(self):
+        return self.lm_head
+
+    def set_output_embeddings(self, new_embeddings):
+        self.lm_head = new_embeddings
+
+    def prepare_inputs_for_generation(self, input_ids, past_key_values=None, inputs_embeds=None, **kwargs):
+        token_type_ids = kwargs.get("token_type_ids", None)
+        # only last token for inputs_ids if past is defined in kwargs
+        if past_key_values:
+            input_ids = input_ids[:, -1].unsqueeze(-1)
+            if token_type_ids is not None:
+                token_type_ids = token_type_ids[:, -1].unsqueeze(-1)
+
+        attention_mask = kwargs.get("attention_mask", None)
+        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[:, -1].unsqueeze(-1)
+        else:
+            position_ids = None
+
+        # 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(
+            {
+                "past_key_values": past_key_values,
+                "use_cache": kwargs.get("use_cache"),
+                "position_ids": position_ids,
+                "attention_mask": attention_mask,
+                "token_type_ids": token_type_ids,
+            }
+        )
+        return model_inputs
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING)
+    @add_code_sample_docstrings(
+        checkpoint=_CHECKPOINT_FOR_DOC,
+        output_type=CausalLMOutputWithCrossAttentions,
+        config_class=_CONFIG_FOR_DOC,
+    )
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Tuple[Tuple[torch.Tensor]]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        encoder_hidden_states: Optional[torch.Tensor] = None,
+        encoder_attention_mask: 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, CausalLMOutputWithCrossAttentions]:
+        r"""
+        labels (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Labels for language modeling. Note that the labels **are shifted** inside the model, i.e. you can set
+            `labels = input_ids` Indices are selected in `[-100, 0, ..., config.vocab_size]` All labels set to `-100`
+            are ignored (masked), the loss is only computed for labels in `[0, ..., config.vocab_size]`
+        """
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        transformer_outputs = self.transformer(
+            input_ids,
+            past_key_values=past_key_values,
+            attention_mask=attention_mask,
+            token_type_ids=token_type_ids,
+            position_ids=position_ids,
+            head_mask=head_mask,
+            inputs_embeds=inputs_embeds,
+            encoder_hidden_states=encoder_hidden_states,
+            encoder_attention_mask=encoder_attention_mask,
+            use_cache=use_cache,
+            output_attentions=output_attentions,
+            output_hidden_states=output_hidden_states,
+            return_dict=return_dict,
+        )
+        hidden_states = transformer_outputs[0]
+
+        # Set device for model parallelism
+        if self.model_parallel:
+            torch.cuda.set_device(self.transformer.first_device)
+            hidden_states = hidden_states.to(self.lm_head.weight.device)
+
+        lm_logits = self.lm_head(hidden_states)
+
+        loss = None
+        if labels is not None:
+            # move labels to correct device to enable model parallelism
+            labels = labels.to(lm_logits.device)
+            # Shift so that tokens < n predict n
+            shift_logits = lm_logits[..., :-1, :].contiguous()
+            shift_labels = labels[..., 1:].contiguous()
+            # Flatten the tokens
+            loss_fct = CrossEntropyLoss()
+            loss = loss_fct(shift_logits.view(-1, shift_logits.size(-1)), shift_labels.view(-1))
+
+        if not return_dict:
+            output = (lm_logits,) + transformer_outputs[1:]
+            return ((loss,) + output) if loss is not None else output
+
+        return CausalLMOutputWithCrossAttentions(
+            loss=loss,
+            logits=lm_logits,
+            past_key_values=transformer_outputs.past_key_values,
+            hidden_states=transformer_outputs.hidden_states,
+            attentions=transformer_outputs.attentions,
+            cross_attentions=transformer_outputs.cross_attentions,
+        )
+
+    @staticmethod
+    def _reorder_cache(
+        past_key_values: Tuple[Tuple[torch.Tensor]], beam_idx: torch.Tensor
+    ) -> Tuple[Tuple[torch.Tensor]]:
+        """
+        This function is used to re-order the `past_key_values` cache if [`~PreTrainedModel.beam_search`] or
+        [`~PreTrainedModel.beam_sample`] is called. This is required to match `past_key_values` with the correct
+        beam_idx at every generation step.
+        """
+        return tuple(
+            tuple(past_state.index_select(0, beam_idx.to(past_state.device)) for past_state in layer_past)
+            for layer_past in past_key_values
+        )
+
+
+@add_start_docstrings(
+    """
+The GPT2 Model transformer with a language modeling and a multiple-choice classification head on top e.g. for
+RocStories/SWAG tasks. The two heads are two linear layers. The language modeling head has its weights tied to the
+input embeddings, the classification head takes as input the input of a specified classification token index in the
+input sequence).
+""",
+    GPT2_START_DOCSTRING,
+)
+class GPT2DoubleHeadsModel(GPT2PreTrainedModel):
+    _tied_weights_keys = ["lm_head.weight"]
+
+    def __init__(self, config):
+        super().__init__(config)
+        config.num_labels = 1
+        self.transformer = GPT2Model(config)
+        self.lm_head = nn.Linear(config.n_embd, config.vocab_size, bias=False)
+        self.multiple_choice_head = SequenceSummary(config)
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    @add_start_docstrings(PARALLELIZE_DOCSTRING)
+    def parallelize(self, device_map=None):
+        warnings.warn(
+            "`GPT2DoubleHeadsModel.parallelize` is deprecated and will be removed in v5 of Transformers, you should"
+            " load your model with `device_map='balanced'` in the call to `from_pretrained`. You can also provide your"
+            " own `device_map` but it needs to be a dictionary module_name to device, so for instance"
+            " {'transformer.h.0': 0, 'transformer.h.1': 1, ...}",
+            FutureWarning,
+        )
+        self.device_map = (
+            get_device_map(len(self.transformer.h), range(torch.cuda.device_count()))
+            if device_map is None
+            else device_map
+        )
+        assert_device_map(self.device_map, len(self.transformer.h))
+        self.transformer.parallelize(self.device_map)
+        self.lm_head = self.lm_head.to(self.transformer.first_device)
+        self.multiple_choice_head = self.multiple_choice_head.to(self.transformer.first_device)
+        self.model_parallel = True
+
+    @add_start_docstrings(DEPARALLELIZE_DOCSTRING)
+    def deparallelize(self):
+        warnings.warn(
+            "Like `parallelize`, `deparallelize` is deprecated and will be removed in v5 of Transformers.",
+            FutureWarning,
+        )
+        self.transformer.deparallelize()
+        self.transformer = self.transformer.to("cpu")
+        self.lm_head = self.lm_head.to("cpu")
+        self.multiple_choice_head = self.multiple_choice_head.to("cpu")
+        self.model_parallel = False
+        torch.cuda.empty_cache()
+
+    def get_output_embeddings(self):
+        return self.lm_head
+
+    def set_output_embeddings(self, new_embeddings):
+        self.lm_head = new_embeddings
+
+    def prepare_inputs_for_generation(self, input_ids, past_key_values=None, **kwargs):
+        token_type_ids = kwargs.get("token_type_ids", None)
+        # only last token for inputs_ids if past is defined in kwargs
+        if past_key_values:
+            input_ids = input_ids[:, -1].unsqueeze(-1)
+            if token_type_ids is not None:
+                token_type_ids = token_type_ids[:, -1].unsqueeze(-1)
+
+        attention_mask = kwargs.get("attention_mask", None)
+        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[:, -1].unsqueeze(-1)
+        else:
+            position_ids = None
+
+        return {
+            "input_ids": input_ids,
+            "past_key_values": past_key_values,
+            "use_cache": kwargs.get("use_cache"),
+            "position_ids": position_ids,
+            "attention_mask": attention_mask,
+            "token_type_ids": token_type_ids,
+        }
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING)
+    @replace_return_docstrings(output_type=GPT2DoubleHeadsModelOutput, config_class=_CONFIG_FOR_DOC)
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Tuple[Tuple[torch.Tensor]]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        mc_token_ids: Optional[torch.LongTensor] = None,
+        labels: Optional[torch.LongTensor] = None,
+        mc_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,
+        **kwargs,
+    ) -> Union[Tuple, GPT2DoubleHeadsModelOutput]:
+        r"""
+        mc_token_ids (`torch.LongTensor` of shape `(batch_size, num_choices)`, *optional*, default to index of the last token of the input):
+            Index of the classification token in each input sequence. Selected in the range `[0, input_ids.size(-1) -
+            1]`.
+        labels (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *optional*):
+            Labels for language modeling. Note that the labels **are shifted** inside the model, i.e. you can set
+            `labels = input_ids`. Indices are selected in `[-100, 0, ..., config.vocab_size - 1]`. All labels set to
+            `-100` are ignored (masked), the loss is only computed for labels in `[0, ..., config.vocab_size - 1]`
+        mc_labels (`torch.LongTensor` of shape `(batch_size)`, *optional*):
+            Labels for computing the multiple choice classification loss. Indices should be in `[0, ..., num_choices]`
+            where *num_choices* is the size of the second dimension of the input tensors. (see *input_ids* above)
+
+        Return:
+
+        Example:
+
+        ```python
+        >>> import torch
+        >>> from transformers import AutoTokenizer, GPT2DoubleHeadsModel
+
+        >>> tokenizer = AutoTokenizer.from_pretrained("gpt2")
+        >>> model = GPT2DoubleHeadsModel.from_pretrained("gpt2")
+
+        >>> # Add a [CLS] to the vocabulary (we should train it also!)
+        >>> num_added_tokens = tokenizer.add_special_tokens({"cls_token": "[CLS]"})
+        >>> # Update the model embeddings with the new vocabulary size
+        >>> embedding_layer = model.resize_token_embeddings(len(tokenizer))
+
+        >>> choices = ["Hello, my dog is cute [CLS]", "Hello, my cat is cute [CLS]"]
+        >>> encoded_choices = [tokenizer.encode(s) for s in choices]
+        >>> cls_token_location = [tokens.index(tokenizer.cls_token_id) for tokens in encoded_choices]
+
+        >>> input_ids = torch.tensor(encoded_choices).unsqueeze(0)  # Batch size: 1, number of choices: 2
+        >>> mc_token_ids = torch.tensor([cls_token_location])  # Batch size: 1
+
+        >>> outputs = model(input_ids, mc_token_ids=mc_token_ids)
+        >>> lm_logits = outputs.logits
+        >>> mc_logits = outputs.mc_logits
+        ```"""
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        transformer_outputs = self.transformer(
+            input_ids,
+            past_key_values=past_key_values,
+            attention_mask=attention_mask,
+            token_type_ids=token_type_ids,
+            position_ids=position_ids,
+            head_mask=head_mask,
+            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]
+
+        # Set device for model parallelism
+        if self.model_parallel:
+            torch.cuda.set_device(self.transformer.first_device)
+            hidden_states = hidden_states.to(self.lm_head.weight.device)
+
+        lm_logits = self.lm_head(hidden_states)
+        mc_logits = self.multiple_choice_head(hidden_states, mc_token_ids).squeeze(-1)
+
+        mc_loss = None
+        if mc_labels is not None:
+            loss_fct = CrossEntropyLoss()
+            mc_loss = loss_fct(mc_logits.view(-1, mc_logits.size(-1)), mc_labels.view(-1))
+        lm_loss = None
+        if labels is not None:
+            labels = labels.to(lm_logits.device)
+            shift_logits = lm_logits[..., :-1, :].contiguous()
+            shift_labels = labels[..., 1:].contiguous()
+            loss_fct = CrossEntropyLoss()
+            lm_loss = loss_fct(shift_logits.view(-1, shift_logits.size(-1)), shift_labels.view(-1))
+
+        if not return_dict:
+            output = (lm_logits, mc_logits) + transformer_outputs[1:]
+            if mc_loss is not None:
+                output = (mc_loss,) + output
+            return ((lm_loss,) + output) if lm_loss is not None else output
+
+        return GPT2DoubleHeadsModelOutput(
+            loss=lm_loss,
+            mc_loss=mc_loss,
+            logits=lm_logits,
+            mc_logits=mc_logits,
+            past_key_values=transformer_outputs.past_key_values,
+            hidden_states=transformer_outputs.hidden_states,
+            attentions=transformer_outputs.attentions,
+        )
+
+    @staticmethod
+    def _reorder_cache(
+        past_key_values: Tuple[Tuple[torch.Tensor]], beam_idx: torch.Tensor
+    ) -> Tuple[Tuple[torch.Tensor]]:
+        """
+        This function is used to re-order the `past_key_values` cache if [`~PreTrainedModel.beam_search`] or
+        [`~PreTrainedModel.beam_sample`] is called. This is required to match `past_key_values` with the correct
+        beam_idx at every generation step.
+        """
+        return tuple(
+            tuple(past_state.index_select(0, beam_idx.to(past_state.device)) for past_state in layer_past)
+            for layer_past in past_key_values
+        )
+
+
+@add_start_docstrings(
+    """
+    The GPT2 Model transformer with a sequence classification head on top (linear layer).
+
+    [`GPT2ForSequenceClassification`] uses the last token in order to do the classification, as other causal models
+    (e.g. GPT-1) 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).
+    """,
+    GPT2_START_DOCSTRING,
+)
+class GPT2ForSequenceClassification(GPT2PreTrainedModel):
+    def __init__(self, config):
+        super().__init__(config)
+        self.num_labels = config.num_labels
+        self.transformer = GPT2Model(config)
+        self.score = nn.Linear(config.n_embd, self.num_labels, bias=False)
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING)
+    @add_code_sample_docstrings(
+        checkpoint="microsoft/DialogRPT-updown",
+        output_type=SequenceClassifierOutputWithPast,
+        config_class=_CONFIG_FOR_DOC,
+    )
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Tuple[Tuple[torch.Tensor]]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[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.transformer(
+            input_ids,
+            past_key_values=past_key_values,
+            attention_mask=attention_mask,
+            token_type_ids=token_type_ids,
+            position_ids=position_ids,
+            head_mask=head_mask,
+            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, sequence_length = input_ids.shape[:2]
+        else:
+            batch_size, sequence_length = inputs_embeds.shape[:2]
+
+        assert (
+            self.config.pad_token_id is not None or batch_size == 1
+        ), "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:
+                sequence_lengths = (torch.eq(input_ids, self.config.pad_token_id).long().argmax(-1) - 1).to(
+                    logits.device
+                )
+            else:
+                sequence_lengths = -1
+                logger.warning(
+                    f"{self.__class__.__name__} will not detect padding tokens in `inputs_embeds`. Results may be "
+                    "unexpected if using padding tokens in conjunction with `inputs_embeds.`"
+                )
+
+        pooled_logits = logits[torch.arange(batch_size, device=logits.device), sequence_lengths]
+
+        loss = None
+        if labels is not None:
+            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,
+        )
+
+
+@add_start_docstrings(
+    """
+    GPT2 Model with a token classification head on top (a linear layer on top of the hidden-states output) e.g. for
+    Named-Entity-Recognition (NER) tasks.
+    """,
+    GPT2_START_DOCSTRING,
+)
+class GPT2ForTokenClassification(GPT2PreTrainedModel):
+    def __init__(self, config):
+        super().__init__(config)
+        self.num_labels = config.num_labels
+
+        self.transformer = GPT2Model(config)
+        if hasattr(config, "classifier_dropout") and config.classifier_dropout is not None:
+            classifier_dropout = config.classifier_dropout
+        elif hasattr(config, "hidden_dropout") and config.hidden_dropout is not None:
+            classifier_dropout = config.hidden_dropout
+        else:
+            classifier_dropout = 0.1
+        self.dropout = nn.Dropout(classifier_dropout)
+        self.classifier = nn.Linear(config.hidden_size, config.num_labels)
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING)
+    # fmt: off
+    @add_code_sample_docstrings(
+        checkpoint="brad1141/gpt2-finetuned-comp2",
+        output_type=TokenClassifierOutput,
+        config_class=_CONFIG_FOR_DOC,
+        expected_loss=0.25,
+        expected_output=["Lead", "Lead", "Lead", "Position", "Lead", "Lead", "Lead", "Lead", "Lead", "Lead", "Lead", "Lead"],
+    )
+    # fmt: on
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        past_key_values: Optional[Tuple[Tuple[torch.Tensor]]] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[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, TokenClassifierOutput]:
+        r"""
+        labels (`torch.LongTensor` of shape `(batch_size, sequence_length)`, *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.transformer(
+            input_ids,
+            past_key_values=past_key_values,
+            attention_mask=attention_mask,
+            token_type_ids=token_type_ids,
+            position_ids=position_ids,
+            head_mask=head_mask,
+            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]
+        hidden_states = self.dropout(hidden_states)
+        logits = self.classifier(hidden_states)
+
+        loss = None
+        if labels is not None:
+            labels = labels.to(logits.device)
+            loss_fct = CrossEntropyLoss()
+            loss = loss_fct(logits.view(-1, self.num_labels), labels.view(-1))
+
+        if not return_dict:
+            output = (logits,) + transformer_outputs[2:]
+            return ((loss,) + output) if loss is not None else output
+
+        return TokenClassifierOutput(
+            loss=loss,
+            logits=logits,
+            hidden_states=transformer_outputs.hidden_states,
+            attentions=transformer_outputs.attentions,
+        )
+
+
+@add_start_docstrings(
+    """
+    The GPT-2 Model transformer with a span classification head on top for extractive question-answering tasks like
+    SQuAD (a linear layer on top of the hidden-states output to compute `span start logits` and `span end logits`).
+    """,
+    GPT2_START_DOCSTRING,
+)
+class GPT2ForQuestionAnswering(GPT2PreTrainedModel):
+    def __init__(self, config):
+        super().__init__(config)
+        self.num_labels = config.num_labels
+        self.transformer = GPT2Model(config)
+        self.qa_outputs = nn.Linear(config.hidden_size, 2)
+
+        # Model parallel
+        self.model_parallel = False
+        self.device_map = None
+        self.gradient_checkpointing = False
+
+        # Initialize weights and apply final processing
+        self.post_init()
+
+    @add_start_docstrings_to_model_forward(GPT2_INPUTS_DOCSTRING.format("batch_size, sequence_length"))
+    @add_code_sample_docstrings(
+        checkpoint=_CHECKPOINT_FOR_DOC,
+        output_type=QuestionAnsweringModelOutput,
+        config_class=_CONFIG_FOR_DOC,
+        real_checkpoint=_CHECKPOINT_FOR_DOC,
+    )
+    def forward(
+        self,
+        input_ids: Optional[torch.LongTensor] = None,
+        attention_mask: Optional[torch.FloatTensor] = None,
+        token_type_ids: Optional[torch.LongTensor] = None,
+        position_ids: Optional[torch.LongTensor] = None,
+        head_mask: Optional[torch.FloatTensor] = None,
+        inputs_embeds: Optional[torch.FloatTensor] = None,
+        start_positions: Optional[torch.LongTensor] = None,
+        end_positions: Optional[torch.LongTensor] = None,
+        output_attentions: Optional[bool] = None,
+        output_hidden_states: Optional[bool] = None,
+        return_dict: Optional[bool] = None,
+    ) -> Union[Tuple, QuestionAnsweringModelOutput]:
+        r"""
+        start_positions (`torch.LongTensor` of shape `(batch_size,)`, *optional*):
+            Labels for position (index) of the start of the labelled span for computing the token classification loss.
+            Positions are clamped to the length of the sequence (`sequence_length`). Position outside of the sequence
+            are not taken into account for computing the loss.
+        end_positions (`torch.LongTensor` of shape `(batch_size,)`, *optional*):
+            Labels for position (index) of the end of the labelled span for computing the token classification loss.
+            Positions are clamped to the length of the sequence (`sequence_length`). Position outside of the sequence
+            are not taken into account for computing the loss.
+        """
+        return_dict = return_dict if return_dict is not None else self.config.use_return_dict
+
+        outputs = self.transformer(
+            input_ids,
+            attention_mask=attention_mask,
+            token_type_ids=token_type_ids,
+            position_ids=position_ids,
+            head_mask=head_mask,
+            inputs_embeds=inputs_embeds,
+            output_attentions=output_attentions,
+            output_hidden_states=output_hidden_states,
+            return_dict=return_dict,
+        )
+
+        sequence_output = outputs[0]
+
+        logits = self.qa_outputs(sequence_output)
+        start_logits, end_logits = logits.split(1, dim=-1)
+        start_logits = start_logits.squeeze(-1).contiguous()
+        end_logits = end_logits.squeeze(-1).contiguous()
+
+        total_loss = None
+        if start_positions is not None and end_positions is not None:
+            # If we are on multi-GPU, split add a dimension
+            if len(start_positions.size()) > 1:
+                start_positions = start_positions.squeeze(-1).to(start_logits.device)
+            if len(end_positions.size()) > 1:
+                end_positions = end_positions.squeeze(-1).to(end_logits.device)
+            # sometimes the start/end positions are outside our model inputs, we ignore these terms
+            ignored_index = start_logits.size(1)
+            start_positions = start_positions.clamp(0, ignored_index)
+            end_positions = end_positions.clamp(0, ignored_index)
+
+            loss_fct = CrossEntropyLoss(ignore_index=ignored_index)
+            start_loss = loss_fct(start_logits, start_positions)
+            end_loss = loss_fct(end_logits, end_positions)
+            total_loss = (start_loss + end_loss) / 2
+
+        if not return_dict:
+            output = (start_logits, end_logits) + outputs[2:]
+            return ((total_loss,) + output) if total_loss is not None else output
+
+        return QuestionAnsweringModelOutput(
+            loss=total_loss,
+            start_logits=start_logits,
+            end_logits=end_logits,
+            hidden_states=outputs.hidden_states,
+            attentions=outputs.attentions,
+        )

gpst/math_util.py

@@ -0,0 +1,37 @@
+import numpy as np
+import torch
+import torch.nn.functional as F
+
+
+def softmax(logits):
+    max = np.max(logits)
+    logits = logits - max
+    exp_x = np.exp(logits)
+    softmax_x = exp_x / np.sum(exp_x)
+    return softmax_x
+
+
+def max_neg_value(dtype):
+    return -torch.finfo(dtype).max
+
+
+def gumbel_softmax(logits, temperature=1, hard=True, train=False):
+    """
+    ST-gumple-softmax
+    input: [*, seq_len, seq_len]
+    return: flatten --> [*, seq_len, seq_len] an one-hot vector
+    """
+    if train:
+        return F.gumbel_softmax(logits, tau=temperature, hard=hard, dim=-1)
+    else:
+        y = F.softmax(logits, dim=-1)
+        if not hard:
+            return y
+
+        shape = y.size()
+        _, ind = y.max(dim=-1)
+        y_hard = torch.zeros_like(y).view(-1, shape[-1])
+        y_hard.scatter_(1, ind.view(-1, 1), 1)
+        y_hard = y_hard.view(*shape)
+        y_hard = (y_hard - y).detach() + y
+        return y_hard

gpst/r2d2_base.py

@@ -0,0 +1,41 @@
+from torch import nn
+import torch
+
+
+class R2D2Base(nn.Module):
+    def __init__(self, config, **kwargs):
+        super().__init__()
+        self.vocab_size = config.vocab_size
+        self.input_dim = config.hidden_size
+        self.hidden_dim = config.intermediate_size
+        self.window_size = config.window_size
+
+        self.tie_decoder = getattr(config, 'tie_decoder', True)
+        self.cls_token_id = config.cls_token_id
+        self.mask_token_id = config.mask_token_id
+        self.pad_token_id = config.pad_token_id
+        self.bos_token_id = config.bos_token_id
+        self.eos_token_id = config.eos_token_id
+        self.nsp_token_id = config.nsp_token_id
+        self.sum_token_id = config.sum_token_id
+
+    @property
+    def device(self):
+        return next(self.parameters()).device
+
+    @property
+    def eos_vec(self):
+        return self.embedding(torch.tensor([self.eos_token_id]).to(self.device)).squeeze(0)
+
+    @property
+    def bos_vec(self):
+        return self.embedding(torch.tensor([self.bos_token_id]).to(self.device)).squeeze(0)
+
+    def from_pretrain(self, model_path, strict=True):
+        state_dict = torch.load(model_path, map_location=lambda a, b: a)
+        transfered_state_dict = {}
+        for k, v in state_dict.items():
+            new_k = k.replace('module.', '')
+            transfered_state_dict[new_k] = v
+        self.load_state_dict(transfered_state_dict, strict=strict)
+        self._tie_weights()

gpst/r2d2_common.py

@@ -0,0 +1,41 @@
+from collections import namedtuple
+
+
+class CacheSlots:
+    E_IJ = 0
+    LOG_P_IJ_SUM = 1
+    NT_SCORE = 2
+
+
+LMLossParam = namedtuple(
+    'LMLossParam',
+    [
+        'model',
+        'chart_tables',
+        'tensor_cache',
+        'input_ids',
+        'flatten_input_ids',
+        's_indices',
+        'atom_spans',
+        'seq_lens'
+    ]
+)
+
+NodeCombination = namedtuple(
+    'NodeCombination',
+    [
+        'node_ik',
+        'node_kj',
+        'left',
+        'right'
+    ]
+)
+
+BOS_CACHE_ID = 0
+EOS_CACHE_ID = 1
+INF_LOG_P_ID = 2
+SPECIAL_TOKEN_NUM = 3  # BOS, EOS
+
+ROLE_LEFT = 1
+ROLE_RIGHT = 2
+ROLE_PARENT = 3

gpst/r2d2_insideoutside.py

@@ -0,0 +1,462 @@
+# coding=utf-8
+# Copyright (c) 2024 Ant Group
+# Author: Xiang Hu
+from typing import List
+import torch.nn as nn
+from typing import Optional
+from gpst.data_structure.py_backend import CPPChartTableManager
+from gpst.fast_parser import TransformerParser
+from gpst.tree_encoder import InsideEncoder, OutsideEncoder
+import torch
+from torch.utils.checkpoint import checkpoint
+from gpst.data_structure.tensor_cache import TensorCache, CacheType
+import torch.nn.functional as F
+from datetime import datetime
+from gpst.math_util import gumbel_softmax
+from dataclasses import dataclass
+
+SPECIAL_TOKEN_NUM = 3
+
+class WeightedSumFunc(torch.autograd.Function):
+    a_ij_require_grad=True
+
+    @staticmethod
+    def forward(ctx, a_ij, e_ijk):
+        # a_ij (N, group_size) e_ijk (N, group_size, dim)
+        ctx.save_for_backward(a_ij, e_ijk)
+        output = torch.einsum('bi,bij->bj', a_ij, e_ijk)
+        return output
+
+    @staticmethod
+    def backward(ctx, grad_output):
+        # (N, dim)
+        a_ij, e_ijk = ctx.saved_tensors
+        if WeightedSumFunc.a_ij_require_grad:
+            d_a_ij = torch.einsum('bj,bij->bi', grad_output.float(), e_ijk)
+        else:
+            d_a_ij = None
+        d_e_ijk = torch.einsum('bi,bj->bij', a_ij, grad_output.float())
+        return d_a_ij, d_e_ijk
+
+@dataclass
+class InsideOutsideContext:
+    scores: Optional = None, 
+    attention_mask: Optional = None, 
+    split_masks: Optional = None, 
+    split_points: Optional = None,
+    batch_size: Optional = None, 
+    root_ids: Optional = None, 
+    inside_cache: Optional = None, 
+    outside_groups: Optional = None,
+    input_cache_ids: Optional = None
+
+DEFAULT_HEIGHT_THRESHOLD=15
+
+class R2D2Base(nn.Module):
+    def __init__(self, config, **kwargs):
+        super().__init__()
+        self.vocab_size = config.vocab_size
+        self.input_dim = config.hidden_size
+        self.hidden_dim = config.intermediate_size
+        self.window_size = config.window_size
+
+        self.tie_decoder = getattr(config, 'tie_decoder', True)
+        self.cls_token_id = config.cls_token_id
+        self.mask_token_id = config.mask_token_id
+        self.pad_token_id = config.pad_token_id
+        self.bos_token_id = config.bos_token_id
+        self.eos_token_id = config.eos_token_id
+        self.nsp_token_id = config.nsp_token_id
+        self.sum_token_id = config.sum_token_id
+
+    @property
+    def device(self):
+        return next(self.parameters()).device
+
+    @property
+    def eos_vec(self):
+        return self.embedding(torch.tensor([self.eos_token_id]).to(self.device)).squeeze(0)
+
+    @property
+    def bos_vec(self):
+        return self.embedding(torch.tensor([self.bos_token_id]).to(self.device)).squeeze(0)
+
+    def from_pretrain(self, model_path, strict=True):
+        state_dict = torch.load(model_path, map_location=lambda a, b: a)
+        transfered_state_dict = {}
+        for k, v in state_dict.items():
+            new_k = k.replace('module.', '')
+            transfered_state_dict[new_k] = v
+        self.load_state_dict(transfered_state_dict, strict=strict)
+        self._tie_weights()
+
+class InsideOutsideModule(R2D2Base):
+    def __init__(self, config, **kwargs):
+        super().__init__(config)
+        self.parser_chunked = config.parser_chunked
+
+        self.parser = TransformerParser(config)
+
+        self.inside_enc = InsideEncoder(config)
+        self.outside_enc = OutsideEncoder(config)
+        self.outside_root_embedding = nn.Parameter(torch.rand(config.hidden_size))
+        self.norm = nn.InstanceNorm1d(config.hidden_size)
+
+        if config.ext_vocab_size > 0:
+            self.ext_embeds = nn.Embedding(config.ext_vocab_size + 1, self.input_dim, padding_idx=0)
+            # initialize with zero
+            self.ext_embeds.weight.data.fill_(0.0)
+
+        self.height_threshold = DEFAULT_HEIGHT_THRESHOLD
+        if hasattr(config, 'height_threshold'):
+            self.height_threshold = config.height_threshold
+        self.use_gumbel = False
+        if hasattr(config, 'use_gumbel'):
+            self.use_gumbel = config.use_gumbel
+        self.ldr_detach = False
+        if hasattr(config, 'ldr_detach'):
+            self.ldr_detach = config.ldr_detach
+
+        self.e_ij_id = -1
+        self.score_sum_id = -1
+        self.score_ijk = -1
+        self.height_ij = -1
+
+        self.reduce_id = config.reduce_token_id
+            
+    def create_tensor_cache(self, seq_lens, total_cache_size=-1):
+        # e_ij, log_p_ij, log_p_sum_ij
+        tensor_cache = TensorCache(
+            self.window_size,
+            seq_lens,
+            cache_types=[
+                CacheType.NORMAL, CacheType.DETACH, 
+                CacheType.NORMAL, CacheType.NORMAL
+            ],
+            dims=[self.input_dim, 1, 1, 1],
+            placeholder_num=SPECIAL_TOKEN_NUM,
+            device=self.device,
+            total_cache_size=total_cache_size)
+        self.e_ij_id = 0
+        self.score_sum_id = 1
+        self.score_ijk = 2
+        self.height_ij = 3
+        tensor_cache.fill(0, tensor_cache.capacity, [self.height_ij], [0])
+        return tensor_cache
+    
+    def _flatten_inputs(self, input_ids, seq_lens, r2d2_embeddings):
+        # Initialize embeddings
+        block_size = input_ids.shape[-1]
+        indices_gather = []
+        for seq_i, seq_len in enumerate(seq_lens):
+            indices_gather.extend(
+                range(block_size * seq_i, block_size * seq_i + seq_len))
+
+        flatten_input_ids = input_ids.flatten()
+        indices_gather = torch.tensor(indices_gather, device=self.device)
+        flatten_input_ids = flatten_input_ids.gather(
+            dim=0, index=indices_gather)
+        flatten_r2d2_emb = r2d2_embeddings.view(-1, r2d2_embeddings.shape[-1]).gather(
+            dim=0, index=indices_gather.unsqueeze(1).repeat(1, r2d2_embeddings.shape[-1])
+        )
+            
+        return flatten_input_ids, flatten_r2d2_emb
+
+    def prepare_composition(self, group_ids, log_p_ids, tensor_cache):
+        e_ij, h_ij = tensor_cache.gather(group_ids.flatten(), [self.e_ij_id, self.height_ij])
+        log_p_ij = tensor_cache.gather(log_p_ids.flatten(), [self.score_sum_id])[0]
+        e_ij = e_ij.view(*group_ids.shape, self.input_dim)
+        h_ij = h_ij.view(*group_ids.shape)  # (batch_size, group_size, 2)
+        log_p_ij = log_p_ij.view(*group_ids.shape) # (batch_size, group_size, 2)
+
+        return e_ij, log_p_ij.sum(dim=-1), h_ij
+
+    def inside(self,
+               inside_cache,
+               span_embeds,
+               temperature,
+               inside_groups):
+        score_orders = []
+        # a_ij_orders = []
+        
+        prepare_time = None
+        inside_time = None
+        weighted_time = None
+        arg_sort_time = None
+        for target_cache_ids, span_ids, cache_ids, detach_cache_ids in inside_groups:
+            # target_cache_ids: (?)
+            # cache_ids: (?, group_size, 2)
+            # detach_cache_ids: (?, group_size, 2)
+
+            # if candidate e_ij and log_p is not empty, apply composition function
+            e_ij, scores_ij_sum, h_ij = self.prepare_composition(
+                cache_ids, detach_cache_ids, inside_cache)
+            # # e_ij: (batch_size, group_size, 2, dim), c_ij: (batch_size, 2, dim)
+            
+            if span_embeds is None:
+                scores_ijk, c_ijk = self.inside_enc(e_ij)
+            else:
+                scores_ijk, c_ijk = self.inside_enc(e_ij, span_embeds[span_ids, :])
+            # scores_ijk, c_ijk = checkpoint(self.inside_enc, e_ij, use_reentrant=False)
+
+            # expected output put c_ijk: (batch_size, group_size, dim)
+            # log_p_ijk: (batch_size, group_size)
+            # print(scores_ijk.shape)
+            # print(scores_ij_sum.shape)
+            scores_ijk_sum = scores_ijk  # (batch_size, combination_size)
+
+            # assert not torch.any(torch.isinf(log_p_ij_step))
+            if not self.use_gumbel:
+                a_ij = F.softmax(scores_ijk_sum / temperature, dim=-1)
+            else:
+                a_ij = gumbel_softmax(scores_ijk_sum, temperature)
+
+            # (batch_size, combination_size)
+            
+            # c_ij = torch.einsum("ij,ijk->ik", a_ij, c_ijk)
+            c_ij = WeightedSumFunc.apply(a_ij, c_ijk)
+            c_ij = self.norm(c_ij)
+            
+            # c_ij_detach = torch.einsum("ij,ijk->ik", a_ij.detach(), c_ijk)
+            h_ij_next, _ = h_ij.max(dim=-1)  # (batch_size, group_size)
+            h_ij_next = h_ij_next + 1
+            h_ij = torch.einsum("ij, ij->i", a_ij, h_ij_next)  # (batch_size)
+            
+            scores_ij_sum = torch.einsum("ij, ij->i", a_ij, scores_ijk_sum).unsqueeze(1)
+
+            inside_cache.scatter(target_cache_ids, [self.e_ij_id, self.score_sum_id, self.height_ij],
+                                [c_ij, scores_ij_sum, h_ij.unsqueeze(1)])
+            
+            # padding to group_size
+            score_orders.append(scores_ijk_sum.argsort(dim=1, descending=True).to('cpu', non_blocking=True))
+            # a_ij_orders.append(a_ij.argsort(dim=1, descending=True).to('cpu', non_blocking=True))
+
+        return score_orders #, a_ij_orders
+
+    def outside_embeddings(self, ctx):
+        root_embedding = self.outside_root_embedding.unsqueeze(0).repeat(ctx.batch_size, 1)
+        outside_cache = self.outside(ctx.batch_size, ctx.root_ids, root_embedding, \
+                                     ctx.inside_cache, ctx.outside_groups)
+        outside_repr = outside_cache.gather(ctx.input_cache_ids, [self.e_ij_id])[0]
+        return outside_repr
+        
+    def parser_loss(self, ctx):
+        # split_masks: (batch_size, L - 1, L - 1)
+        # split points: (batch_size, L - 1)
+        scores = ctx.scores.clone()
+        split_masks = ctx.split_masks.clone()
+        split_points = ctx.split_points.clone()
+        L = scores.shape[1]
+        attention_mask = ctx.attention_mask.clone()
+
+        proxyminusinf = -1e9
+        assert len(attention_mask.shape) == 2
+        scores.masked_fill_(attention_mask[:, 1: L + 1] == 0, proxyminusinf)
+        scores = scores.unsqueeze(1).repeat(1, L, 1)
+        scores.masked_fill_(split_masks[:, :L, :L] == 0, proxyminusinf)  # (batch_size, L - 1, L - 1)
+        #print("Split points: ", split_points.shape, "\n", split_points)
+        # test only feedback on root split
+        # log_p = F.log_softmax(scores.float(), dim=-1)  # (batch_size, L - 1, L - 1)
+        #print("scores stats: ", scores.min().item(), scores.max().item(), scores.mean().item())
+        #print("split_points shape:", split_points.shape)
+        #print("split_points unique:", split_points.unique())
+        #print("split_points max:", split_points.max().item(), "vocab size:", scores.size(-1))
+
+        return F.cross_entropy(scores.clone().transpose(1, 2).float(), split_points.clone()[:, :L], ignore_index=-1)  
+
+    def outside(self, batch_size, root_ids, root_embedding, inside_cache, outside_groups):
+        # initialize tensor cache for outside algorithm
+        out_cache_size = inside_cache.capacity - inside_cache.placeholder_num
+        outside_cache = TensorCache(0, None, [CacheType.NORMAL, CacheType.NORMAL, CacheType.NORMAL],
+                                    [self.input_dim, 1, 1], inside_cache.placeholder_num,
+                                    total_cache_size=out_cache_size, 
+                                    device=inside_cache.device)
+        topdown_e_ij_slot = 0
+        topdown_score_slot = 1  # weighted sum for outside scores
+        topdown_score_ln_sum = 2  # store log (e^w1 + e^w2 + e^w3), w1, w2, w3 is the calculated outside scores
+        
+        # (batch_size, dim), add root role embedding
+        
+        zero_padding = torch.zeros(batch_size, 1, dtype=torch.float, device=self.device)
+        neg_padding = torch.zeros((outside_cache.capacity, 1), dtype=torch.float, device=self.device).fill_(-1e20)
+        
+        # As there is no calcuated outside scores, initialize caches with a huge neg value
+        outside_cache.fill(0, outside_cache.capacity, [topdown_score_ln_sum], [neg_padding])
+        outside_cache.scatter(root_ids.long(), [topdown_e_ij_slot, topdown_score_slot, topdown_score_ln_sum], 
+                              [root_embedding, zero_padding, zero_padding])
+
+        # run outside according to inside groups
+        for target_cache_ids, cache_ids, detach_cache_ids in outside_groups:
+            parent_ids = target_cache_ids
+            child_ids = cache_ids # (N, comb_size, 2)
+            
+            # assert child_ids[:, :, 0].unique().shape[0] == cache_ids.shape[0] * cache_ids.shape[1]
+            # assert child_ids[:, :, 1].unique().shape[0] == cache_ids.shape[0] * cache_ids.shape[1]
+
+            score_ids = detach_cache_ids
+            
+            parent_ij, parent_ij_score = outside_cache.gather(parent_ids, [topdown_e_ij_slot, topdown_score_slot])
+
+            child_ids_shape = child_ids.shape  # (batch_size, comb_size, 2)
+            child_ikj = inside_cache.gather(child_ids.flatten(), [self.e_ij_id])[0]
+            child_scores = inside_cache.gather(score_ids.flatten(), [self.score_sum_id])[0]
+            child_ikj = child_ikj.view(*child_ids.shape, -1)
+            child_scores = child_scores.view(*child_ids.shape)  # (batch_size, comb_size, 2)
+
+            out_scores, out_ikj = self.outside_enc(parent_ij, child_ikj, parent_ij_score, child_scores)
+            # out_scores, out_ikj = checkpoint(self.outside_enc, parent_ij, child_ikj, parent_ij_score, child_scores, use_reentrant=False)
+            # span_norm = (1 + max_lens - span_lens).unsqueeze(1).unsqueeze(2)
+            # out_ikj: (batch_size, comb_size, 2, dim)
+            
+            dim = out_ikj.shape[-1]
+
+            weighted_e_ij, weighted_scores, log_ksum_score = \
+                outside_cache.gather(child_ids[:, :, 0].flatten(), 
+                                     [topdown_e_ij_slot, topdown_score_slot, topdown_score_ln_sum])
+            weighted_e_ij = weighted_e_ij.view(*child_ids_shape[:-1], dim)  # (batch_size, comb_size, dim)
+            log_ksum_score = log_ksum_score.view(*child_ids_shape[:-1])  # (batch_size, comb_size)
+            weighted_scores = weighted_scores.view(*child_ids_shape[:-1])
+
+            # log_p_ijk_mean: (batch_size, comb_size)
+            left_k_sum_scores = torch.stack([log_ksum_score, out_scores[:, :, 0]], dim=2)  # (batch_size, comb_size, 2)
+            left_k_weights = F.softmax(left_k_sum_scores, dim=2)
+            left_weighted_e_ij = left_k_weights[:, :, 0].unsqueeze(2) * weighted_e_ij + \
+                                    left_k_weights[:, :, 1].unsqueeze(2) * out_ikj[:, :, 0, :]
+            left_weighted_scores = left_k_weights[:, :, 0] * weighted_scores + \
+                                    left_k_weights[:, :, 1] * out_scores[:, :, 0]
+
+            # (batch_size, comb_size, dim)
+            left_k_sum_scores = left_k_sum_scores.logsumexp(dim=2, keepdim=True)
+
+            left_weighted_e_ij = left_weighted_e_ij.view(-1, dim)
+            left_weighted_scores = left_weighted_scores.view(-1, 1)
+            left_k_sum_scores = left_k_sum_scores.view(-1, 1)
+
+            outside_cache.scatter(child_ids[:, :, 0].flatten().long(), 
+                                  [topdown_e_ij_slot, topdown_score_slot, topdown_score_ln_sum], 
+                                  [left_weighted_e_ij, left_weighted_scores, left_k_sum_scores])
+            
+            weighted_e_ij, weighted_scores, log_ksum_score = \
+                outside_cache.gather(child_ids[:, :, 1].flatten(), 
+                                     [topdown_e_ij_slot, topdown_score_slot, topdown_score_ln_sum])
+            weighted_e_ij = weighted_e_ij.view(*child_ids_shape[:-1], dim)  # (batch_size, comb_size, dim)
+            log_ksum_score = log_ksum_score.view(*child_ids_shape[:-1])  # (batch_size, comb_size)
+            weighted_scores = weighted_scores.view(*child_ids_shape[:-1])
+
+            right_k_sum_scores = torch.stack([log_ksum_score, out_scores[:, :, 1]], dim=2)  # (batch_size, comb_size, 2)
+            right_k_weights = F.softmax(right_k_sum_scores, dim=2)
+            right_weighted_e_ij = right_k_weights[:, :, 0].unsqueeze(2) * weighted_e_ij + \
+                                    right_k_weights[:, :, 1].unsqueeze(2) * out_ikj[:, :, 1, :]
+            right_weighted_scores = right_k_weights[:, :, 0] * weighted_scores + \
+                                    right_k_weights[:, :, 1] * out_scores[:, :, 1]
+
+            # (batch_size, comb_size, dim)
+            right_k_sum_scores = right_k_sum_scores.logsumexp(dim=2, keepdim=True)
+
+            right_weighted_e_ij = right_weighted_e_ij.view(-1, dim)
+            right_weighted_scores = right_weighted_scores.view(-1, 1)
+            right_k_sum_scores = right_k_sum_scores.view(-1, 1)
+            
+            outside_cache.scatter(child_ids[:, :, 1].flatten().long(), 
+                                  [topdown_e_ij_slot, topdown_score_slot, topdown_score_ln_sum], 
+                                  [right_weighted_e_ij, right_weighted_scores, right_k_sum_scores])
+
+        return outside_cache
+
+    def forward(self, 
+                chunk_input_ids,
+                chunk_masks,
+                input_ids,
+                masks,
+                r2d2_embeddings, # corresponding to chunked_input_ids
+                group_ids,
+                max_input_len,
+                atom_spans:List[List[int]]=None,
+                eos_labels=None,
+                span_ids=None,
+                external_vocab_ids=None,
+                coeff=1.0,
+                temperature=1.0):
+
+        split_indices, split_scores = self.parser(chunk_input_ids, chunk_masks, atom_spans=atom_spans, noise_coeff=coeff)
+        split_indices = split_indices.to('cpu', non_blocking=True)
+
+        seq_lens = torch.sum(masks, dim=1, dtype=torch.int)  # (batch_size)
+        seq_lens_np = seq_lens.to('cpu').data.numpy()
+        
+        if len(chunk_masks.shape) == 2:
+            chunk_seq_lens_np = (chunk_masks != 0).sum(dim=1).cpu().data.numpy()
+            # chunk_seq_lens_np = chunk_masks.sum(dim=1).cpu().data.numpy()
+        # elif len(chunk_masks.shape) == 3:
+        #     chunk_seq_lens_np = (chunk_masks.sum(dim=1) > 0).cpu().to(int).sum(dim=1).data.numpy()
+
+        batch_size = input_ids.shape[0]
+        input_ids_cpu = input_ids.to('cpu', non_blocking=True)
+        
+        flatten_input_ids, flatten_r2d2_emb = self._flatten_inputs(chunk_input_ids, chunk_seq_lens_np, r2d2_embeddings)
+        ids_num = flatten_input_ids.shape[0]
+        input_cache_ids = torch.arange(SPECIAL_TOKEN_NUM, 
+                                       SPECIAL_TOKEN_NUM + ids_num).to(self.device)
+        
+        inside_cache = self.create_tensor_cache(seq_lens_np)
+        inside_cache.scatter(input_cache_ids, [self.e_ij_id], [flatten_r2d2_emb])
+
+        tables = CPPChartTableManager(seq_lens_np, self.window_size, split_indices.data.numpy(),
+                                      inside_cache.placeholder_num, inside_cache.detach_offset, group_ids=group_ids,
+                                      span_ids=span_ids)
+        target_cache_ids, span_ids, cache_ids, detach_cache_ids = \
+                tables.construct_inside_groups(self.device)
+        root_ids = tables.root_ids
+
+        span_embeds = None
+        if external_vocab_ids is not None:
+            span_embeds = self.ext_embeds(external_vocab_ids)
+
+        score_orders = self.inside(inside_cache, span_embeds, temperature,
+                                    zip(target_cache_ids, span_ids, cache_ids, detach_cache_ids))
+        
+        span_masks, split_targets, ldr_cache_ids, position_ids, tgt_ids, token_indices, ext_ids = \
+            tables.prepare_generation(score_orders, score_orders, atom_spans, input_ids_cpu.data.numpy(),
+                                      group_ids, self.eos_token_id, self.reduce_id,
+                                      max_input_len, eos_labels=eos_labels)
+        # span_mask, split_targets, ldr_cache_ids, position_ids, tgt_ids
+
+        ldr_cache_ids = ldr_cache_ids.clone().to(self.device, non_blocking=True)
+        position_ids = position_ids.to(self.device, non_blocking=True)
+        tgt_ids = tgt_ids.to(self.device, non_blocking=True)
+        ext_ids = ext_ids.to(self.device, non_blocking=True)
+        token_indices = token_indices.to(self.device, non_blocking=True)
+        
+        span_masks = span_masks.to(self.device, non_blocking=True)
+        split_targets = split_targets.to(self.device, non_blocking=True)
+
+        ldr_repr = inside_cache.gather(ldr_cache_ids.flatten(), [self.e_ij_id])[0]
+        ldr_repr = ldr_repr.view(*ldr_cache_ids.shape, -1)  # (N, L, dim)
+
+        # l_height = (inside_cache.gather(root_ids, [self.height_ij])[0] / seq_lens).mean()
+        inside_height = inside_cache.gather(root_ids, [self.height_ij])[0]
+
+        inside_height = torch.where(inside_height > self.height_threshold, inside_height - self.height_threshold, 0)
+        height_norm = torch.where(seq_lens > self.height_threshold, seq_lens - self.height_threshold, 1)
+        l_height = (inside_height / height_norm).mean()
+        
+        # fix clones
+        tci = [t.clone() for t in target_cache_ids]
+        ci = [t.clone() for t in cache_ids]
+        dci = [t.clone() for t in detach_cache_ids]
+        ctx = InsideOutsideContext(
+            scores=split_scores, 
+            attention_mask=chunk_masks,
+            split_masks=span_masks.clone(), 
+            split_points=split_targets.clone(),
+            batch_size=batch_size, 
+            root_ids=root_ids, 
+            inside_cache=inside_cache, 
+            input_cache_ids=input_cache_ids,
+            outside_groups=list(zip(reversed(tci), reversed(ci), reversed(dci)))
+        )
+        if self.ldr_detach:
+            ldr_repr = ldr_repr.detach()
+        
+        return ctx, flatten_input_ids, ldr_repr, position_ids.clone(), \
+            tgt_ids.clone(), token_indices.clone(), ext_ids, split_targets.clone(), l_height

gpst/setup.py

@@ -0,0 +1,20 @@
+from setuptools import setup
+import os
+import torch
+import glob
+from torch.utils.cpp_extension import BuildExtension, CUDAExtension, CppExtension
+
+
+include_dirs = os.path.dirname(os.path.abspath(__file__))
+source_file = glob.glob(os.path.join('./', 'cpp_extension', '*.cpp'))
+
+if torch.cuda.is_available():
+    setup(
+        name='cppbackend',
+        ext_modules=[
+            CppExtension('cppbackend',
+                          sources=source_file,
+                          include_dirs=[include_dirs]
+                         )
+        ],
+        cmdclass={'build_ext': BuildExtension}) 

gpst/topdown_parser.py

@@ -0,0 +1,69 @@
+from copy import deepcopy
+from typing import List, Tuple
+import torch.nn.functional as F
+import torch.nn as nn
+import torch
+import numpy as np
+from torch.nn.utils.rnn import pack_padded_sequence, pad_packed_sequence
+from gpst.tree_encoder import _get_activation_fn
+
+
+INF=1e7
+
+class BasicParser(nn.Module):
+    def __init__(self):
+        super().__init__()
+
+    def _split_point_scores(self, input_ids, attn_mask):
+        pass
+
+    def _adjust_atom_span(self, scores, atom_spans, const=1):
+        # scores.masked_fill_(attention_mask[:, 1:scores.shape[1] + 1] == 0, float('-inf'))
+        points_mask = np.full(scores.shape, fill_value=0)
+        for batch_i, spans in enumerate(atom_spans):
+            if spans is not None:
+                for (i, j) in spans:
+                    points_mask[batch_i][i: j] += 1
+        points_mask = torch.tensor(points_mask, device=scores.device)
+        assert const > 0
+        mask_scores = points_mask * (scores.max() - scores.min() + const)
+        return scores - mask_scores
+
+    # @torch.inference_mode
+    def parse(self, input_ids: torch.Tensor = None, attention_mask: torch.Tensor = None,
+              atom_spans: List[List[Tuple[int]]] = None, noise_coeff: float = 1.0):
+        """
+        params:
+            input_ids: torch.Tensor, 
+            attention_mask:
+            atom_spans: List[List[Tuple[int]]], batch_size * span_lens * 2, each span contains start and end position
+            splits: List[List[int]], batch_size * split_num, list of split positions
+        """
+        org_scores = self._split_point_scores(input_ids, attention_mask)
+        # meaningful split points: seq_lens - 1
+
+        if self.training:
+            noise = -torch.empty_like(
+                org_scores,
+                memory_format=torch.legacy_contiguous_format,
+                requires_grad=False).exponential_().log() * max(0, noise_coeff)
+            scores = org_scores.detach() + noise
+        else:
+            scores = org_scores.detach()
+        if atom_spans is not None:
+            scores = self._adjust_atom_span(scores, atom_spans)
+
+        if attention_mask is not None:
+            if len(attention_mask.shape) == 3:
+                attention_mask = (attention_mask.sum(dim=1) > 0).to(int)
+            scores = scores.masked_fill_(attention_mask[:, 1:scores.shape[1] + 1] == 0, float('inf'))
+        # split according to scores
+        # for torch >= 1.9
+        _, s_indices = scores.sort(dim=-1, descending=False, stable=True)
+        return s_indices, org_scores
+
+    def forward(self, input_ids: torch.Tensor = None, attention_mask: torch.Tensor = None,
+                atom_spans: List[List[Tuple[int]]] = None, noise_coeff: float = 1.0):
+        return self.parse(input_ids, attention_mask=attention_mask, atom_spans=atom_spans, 
+                          noise_coeff=noise_coeff)
+

gpst/tree_encoder.py

@@ -0,0 +1,225 @@
+import math
+from typing import List
+import torch.nn as nn
+import torch.nn.functional as F
+import torch
+from torch.nn import init
+from copy import deepcopy
+import numpy as np
+
+ACTIVATION_POOL = ['relu', 'gelu']
+
+
+def _get_activation_fn(activation):
+    if activation in ACTIVATION_POOL:
+        return getattr(F, activation)
+
+    raise RuntimeError("activation should be relu/gelu, not {}".format(activation))
+
+
+class TreeEncoderLayer(nn.Module):
+    def __init__(self, d_model, nhead, dim_feedforward, dropout, max_role_count, 
+                 activation='gelu', batch_first=False, val_position=False):
+        super().__init__()
+        self.self_attn = nn.MultiheadAttention(d_model, nhead, dropout=dropout, batch_first=batch_first)
+        self._val_position = val_position
+        self.linear1 = nn.Linear(d_model, dim_feedforward)
+        self.dropout = nn.Dropout(dropout)
+        self.linear2 = nn.Linear(dim_feedforward, d_model)
+
+        self.norm1 = nn.InstanceNorm1d(d_model)
+        self.norm2 = nn.InstanceNorm1d(d_model)
+        self.dropout1 = nn.Dropout(dropout)
+        self.dropout2 = nn.Dropout(dropout)
+        self.position_embedding = nn.Embedding(max_role_count, d_model)
+
+        self.activation = _get_activation_fn(activation)
+
+    def forward(self, src, src_mask=None, pos_ids=None):
+        """
+        :param src: concatenation of task embeddings and representation for left and right.
+                    src shape: (task_embeddings + left + right, batch_size, dim)
+        :param src_mask:
+        :param pos_ids:
+        :return:
+        """
+        if len(pos_ids.shape) == 1:
+            sz = src.shape[1]  # sz: batch_size
+            pos_ids = pos_ids.unsqueeze(1).expand(-1, sz)  # (3, batch_size)
+        position_embedding = self.position_embedding(pos_ids)
+        src2 = self.self_attn(src + position_embedding, src + position_embedding, src, 
+                              attn_mask=src_mask)[0]
+        src = src + self.dropout1(src2)
+        src = self.norm1(src)
+        src2 = self.linear2(self.dropout(self.activation(self.linear1(src))))
+        src = src + self.dropout2(src2)
+        src = self.norm2(src)
+        return src
+
+class InsideEncoder(nn.Module):
+    def __init__(self, config):
+        super().__init__()
+
+        self.const_size = config.const_size
+        self.left_linear = nn.Sequential(nn.Linear(config.hidden_size, config.hidden_size),
+                                         nn.GELU(),
+                                         nn.Linear(config.hidden_size, config.const_size))
+        self.right_linear = nn.Sequential(nn.Linear(config.hidden_size, config.hidden_size),
+                                          nn.GELU(),
+                                          nn.Linear(config.hidden_size, config.const_size))
+        # self.const_linear = nn.Sequential(GroupLinear(2, config.hidden_size, config.hidden_size),
+        #                                   nn.GELU(),
+        #                                   GroupLinear(2, config.hidden_size, config.hidden_size))
+        
+        layer = TreeEncoderLayer(config.hidden_size,
+                                 config.num_attention_heads,
+                                 config.intermediate_size,
+                                 max_role_count=config.max_role_embeddings,
+                                 dropout=config.attention_probs_dropout_prob,
+                                 activation='gelu',
+                                 batch_first=True,
+                                 val_position=True)
+        self.norm = nn.InstanceNorm1d(config.hidden_size)
+        self.layers = nn.ModuleList([layer] + [deepcopy(layer) for _ in range(config.encoder_num_hidden_layers - 1)])
+        self._device = None
+        self._pos_ids = None
+        
+    
+    @property
+    def device(self):
+        if self._device is None:
+            self._device = next(self.parameters()).device
+        return self._device
+    
+    @property
+    def pos_ids(self):
+        if self._pos_ids is None:
+            self._pos_ids = torch.arange(2).to(self.device)
+        return self._pos_ids
+
+    def forward(self, src, span_embeds=None):
+        """
+        :param src: [batch_size, comb_size, 2, dim]
+        :param span_embeds: [batch_size, dim]
+        :return:
+        """
+        dim = src.shape[-1]
+        org_shape = src.shape  # (batch_size, comb_size, 2, dim)
+        output = src.view(-1, 2, dim)
+
+        # torch.cuda.synchronize()
+        # with torch.cuda.stream(self.s1):
+        left_const = self.left_linear(output[:, 0, :])
+        right_const = self.right_linear(output[:, 1, :])
+        
+        for mod in self.layers:
+            output = mod(output, pos_ids=self.pos_ids.unsqueeze(0))
+
+        mat_scores = torch.einsum("bi,bi->b", left_const, right_const) / math.sqrt(self.const_size)
+        mat_scores = mat_scores.view(*org_shape[:-2])  # (batch_size, comb_size)
+        if span_embeds is not None:
+            output = output.sum(dim=1).view(*org_shape[:-2], dim) + span_embeds.unsqueeze(1)
+        else:
+            output = output.sum(dim=1).view(*org_shape[:-2], dim)
+
+        return mat_scores, self.norm(output)
+
+    
+class OutsideEncoder(nn.Module):
+    def __init__(self, config) -> None:
+        super().__init__()
+        
+        self.parent_linear = nn.Sequential(nn.Linear(config.hidden_size, config.hidden_size),
+                                           nn.GELU(),
+                                           nn.Linear(config.hidden_size, config.const_size))
+        self.left_linear = nn.Sequential(nn.Linear(config.hidden_size, config.hidden_size),
+                                         nn.GELU(),
+                                         nn.Linear(config.hidden_size, config.const_size))
+        self.right_linear = nn.Sequential(nn.Linear(config.hidden_size, config.hidden_size),
+                                          nn.GELU(),
+                                          nn.Linear(config.hidden_size, config.const_size))
+        # self.const_linear = nn.Sequential(GroupLinear(3, config.hidden_size, config.hidden_size),
+        #                                   nn.GELU(),
+        #                                   GroupLinear(3, config.hidden_size, config.hidden_size))
+        self.const_size = config.const_size
+        
+        layer = TreeEncoderLayer(config.hidden_size,
+                                 config.num_attention_heads,
+                                 config.intermediate_size,
+                                 max_role_count=config.max_role_embeddings,
+                                 dropout=config.attention_probs_dropout_prob,
+                                 activation='gelu',
+                                 batch_first=True,
+                                 val_position=True)
+        self.norm = nn.InstanceNorm1d(config.hidden_size)
+        self.layers = nn.ModuleList([layer] + [deepcopy(layer) for _ in range(config.decoder_num_hidden_layers - 1)])
+        self._device = None
+        self._dec_pos_ids = None
+    
+    @property
+    def device(self):
+        if self._device is None:
+            self._device = next(self.parameters()).device
+        return self._device
+    
+    @property
+    def dec_pos_ids(self):
+        if self._dec_pos_ids is None:
+            self._dec_pos_ids = torch.tensor([[0, 1], [0, 2]], device=self.device)
+        return self._dec_pos_ids 
+    
+    def forward(self, parent_ij, child_ikj, parent_scores=None, child_scores=None):
+        """
+        :param parent_ij: (batch_size, dim)
+        :param parent_scores: (batch_size, 1)
+        :param child_ikj: (batch_size, comb_size, 2, dim)
+        :param child_scores: (batch_size, comb_size, 2)
+        :return: (batch_size, 2), (batch_size, 2, dim)
+        """
+        # p_l = parent_ij @ self.W_outside_r  # (batch_size, dim)
+        # out_score_ik = torch.einsum('bd, bcd->bc', p_l, child_ikj[:, :, 1, :])  # (batch_size, comb_size)
+        # out_score_ik = out_score_ik + parent_scores + child_scores[:, :, 1, 0] # (batch_size, comb_size)
+
+        # p_r = parent_ij @ self.W_outside_l  # (batch_size, dim)
+        # out_score_kj = torch.einsum('bd, bcd->bc', p_r, child_ikj[:, :, 0, :])  # (batch_size, comb_size)
+        # out_score_kj = out_score_kj + parent_scores + child_scores[:, :, 0, 0] # (batch_size, comb_size)
+
+        batch_size = child_ikj.shape[0]
+        comb_size = child_ikj.shape[1]
+        
+        comb_size = child_ikj.shape[1]
+        parent_ij_ext = parent_ij.unsqueeze(1).unsqueeze(2).repeat(1, comb_size, 2, 1)  # (batch_size, comb_size, 2, dim)
+        
+        inputs = torch.stack([parent_ij_ext, child_ikj.flip([2])], dim=3)  # (batch_size, comb_size, 2, 2, dim)
+        inputs = inputs.view(batch_size * comb_size * 2, 2, -1)
+
+
+        outside_scores = None
+        
+        if parent_scores is not None and child_scores is not None:
+            parent_const = self.parent_linear(inputs[:, 0, :])
+            right_child_const = self.right_linear(inputs[::2, 1, :])
+            left_child_const = self.left_linear(inputs[1::2, 1, :])
+
+
+        # self.dec_pos_ids: (2, 2)
+        pos_ids = self.dec_pos_ids.repeat(batch_size * comb_size, 1)
+        for mod in self.layers:
+            inputs = mod(inputs, pos_ids=pos_ids)
+
+        if parent_scores is not None and child_scores is not None:
+            parent_const_r = parent_const[::2, :]
+            parent_const_l = parent_const[1::2, :]
+            left_score = (parent_const_r * right_child_const).sum(dim=-1) / math.sqrt(self.const_size)
+            # right_score = torch.einsum('bi,bi->b', parent_const_l, left_child_const) / math.sqrt(self.const_size)
+            right_score = (parent_const_l * left_child_const).sum(dim=-1) / math.sqrt(self.const_size)
+            left_score = left_score.view(batch_size, comb_size)
+            right_score = right_score.view(batch_size, comb_size)
+            out_score_ik = left_score
+            out_score_kj = right_score
+            outside_scores = torch.stack([out_score_ik, out_score_kj], dim=2)
+        # inputs: (?, 2, dim)
+        out_e_ij = self.norm(inputs.sum(dim=1))
+
+        return outside_scores, \
+               out_e_ij.view(batch_size, comb_size, 2, -1)  # (batch_size, comb_size, 2, dim)