Files
wehub-resource-sync 59a0a3844c
PR Test AMD / cancel-on-close (push) Has been skipped
PR Test NVIDIA ARM / scan (push) Has been skipped
PR Test NVIDIA / cancel-on-close (push) Has been skipped
PR Test AMD / scan (push) Has been skipped
PR Test NVIDIA ARM / cancel-on-close (push) Has been skipped
PR Test NVIDIA / scan (push) Has been skipped
Release Docker Images / build (cu129-torch-2.11.0) (push) Has been skipped
Release Docker Images / build (cu130-torch-2.11.0) (push) Has been skipped
Release PyPI / publish (push) Has been skipped
Scheduler Python Test / test (push) Successful in 27m19s
Docs / build (push) Successful in 28m8s
Scheduler C++ Test / test (push) Successful in 28m19s
Scheduler C++ Test / test-flat (push) Successful in 28m18s
Docs / deploy (push) Has been cancelled
PR Test AMD / finish (push) Has been cancelled
PR Test NVIDIA / finish (push) Has been cancelled
PR Test NVIDIA ARM / finish (push) Has been cancelled
PR Test NVIDIA ARM / ${{ matrix.name }} (${{ matrix.runner }}) (push) Has been cancelled
PR Test AMD / ${{ matrix.name }} (${{ matrix.runner }}) (push) Has been cancelled
PR Test NVIDIA / ${{ matrix.name }} (${{ matrix.runner }}) (push) Has been cancelled
chore: import upstream snapshot with attribution
2026-07-13 12:32:31 +08:00

651 lines
24 KiB
Python

# Copyright (c) 2026 LightSeek Foundation
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in
# all copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
# SOFTWARE.
"""Qwen3 visual tower base reused by the Qwen3.5 target models."""
from __future__ import annotations
from collections.abc import Callable
from functools import lru_cache, partial
import numpy as np
import torch
import torch.nn as nn
from einops import rearrange
from transformers.activations import ACT2FN
from tokenspeed.runtime.configs.qwen3_vision_config import Qwen3VLVisionConfig
from tokenspeed.runtime.distributed.mapping import Mapping
from tokenspeed.runtime.layers.attention.mm_encoder_attention import (
VIT_CUDNN_BATCH_BUCKETS,
VIT_CUDNN_SEQLEN_BUCKETS,
VIT_CUDNN_WORKSPACE_BYTES,
VisionAttention,
round_up_to_bucket,
)
from tokenspeed.runtime.layers.conv import Conv3dLayer
from tokenspeed.runtime.layers.linear import ColumnParallelLinear, RowParallelLinear
from tokenspeed.runtime.layers.quantization.base_config import QuantizationConfig
from tokenspeed.runtime.layers.rotary_embedding import get_rope
from tokenspeed.runtime.layers.vocab_parallel_embedding import VocabParallelEmbedding
from tokenspeed.runtime.utils import add_prefix
@lru_cache(maxsize=1024)
def _rot_pos_ids(h: int, w: int, spatial_merge_size: int) -> torch.Tensor:
if isinstance(h, torch.Tensor):
h = int(h.item())
if isinstance(w, torch.Tensor):
w = int(w.item())
if isinstance(spatial_merge_size, torch.Tensor):
spatial_merge_size = int(spatial_merge_size.item())
hpos_ids = np.broadcast_to(np.arange(h).reshape(h, 1), (h, w))
h_div = h // spatial_merge_size
w_div = w // spatial_merge_size
hpos_ids = hpos_ids.reshape(h_div, spatial_merge_size, w_div, spatial_merge_size)
hpos_ids = hpos_ids.transpose(0, 2, 1, 3).flatten()
wpos_ids = np.broadcast_to(np.arange(w).reshape(1, w), (h, w))
wpos_ids = wpos_ids.reshape(h_div, spatial_merge_size, w_div, spatial_merge_size)
wpos_ids = wpos_ids.transpose(0, 2, 1, 3).flatten()
return torch.from_numpy(np.stack([hpos_ids, wpos_ids], axis=-1))
class Qwen3VLVisionMLP(nn.Module):
def __init__(
self,
in_features: int,
hidden_features: int,
mapping: Mapping,
bias: bool = True,
hidden_act="silu",
quant_config: QuantizationConfig | None = None,
prefix: str = "",
):
super().__init__()
vision = mapping.vision
self.linear_fc1 = ColumnParallelLinear(
in_features,
hidden_features,
bias=bias,
quant_config=quant_config,
prefix=add_prefix("linear_fc1", prefix),
tp_size=vision.tp_size,
tp_rank=vision.tp_rank,
tp_group=vision.tp_group,
)
self.linear_fc2 = RowParallelLinear(
hidden_features,
in_features,
bias=bias,
quant_config=quant_config,
prefix=add_prefix("linear_fc2", prefix),
tp_size=vision.tp_size,
tp_rank=vision.tp_rank,
tp_group=vision.tp_group,
reduce_results=True,
)
self.act = ACT2FN[hidden_act]
def forward(self, x: torch.Tensor):
x_fc1, _ = self.linear_fc1(x)
x_act = self.act(x_fc1)
mlp_output, _ = self.linear_fc2(x_act)
return mlp_output
class Qwen3VLVisionPatchEmbed(nn.Module):
def __init__(self, config) -> None:
super().__init__()
self.patch_size = config.patch_size
self.temporal_patch_size = config.temporal_patch_size
self.in_channels = config.in_channels
self.embed_dim = config.hidden_size
kernel_size = [self.temporal_patch_size, self.patch_size, self.patch_size]
self.proj = Conv3dLayer(
self.in_channels,
self.embed_dim,
kernel_size=kernel_size,
stride=kernel_size,
bias=True,
)
def forward(self, hidden_states: torch.Tensor) -> torch.Tensor:
target_dtype = self.proj.weight.dtype
hidden_states = hidden_states.view(
-1,
self.in_channels,
self.temporal_patch_size,
self.patch_size,
self.patch_size,
)
hidden_states = self.proj(hidden_states.to(dtype=target_dtype)).view(
-1, self.embed_dim
)
return hidden_states
class Qwen3VLVisionBlock(nn.Module):
def __init__(
self,
dim: int,
num_heads: int,
intermediate_dim: int,
mapping: Mapping,
head_size: int | None = None,
hidden_act="silu",
norm_layer: Callable[[int], nn.Module] | None = None,
quant_config: QuantizationConfig | None = None,
prefix: str = "",
workspace_buffer: torch.Tensor | None = None,
mm_attention_backend: str | None = None,
) -> None:
super().__init__()
if norm_layer is None:
norm_layer = partial(nn.LayerNorm, eps=1e-6)
self.norm1 = norm_layer(dim)
self.norm2 = norm_layer(dim)
self.attn = VisionAttention(
embed_dim=dim,
num_heads=num_heads,
head_size=head_size,
proj_bias=True,
quant_config=quant_config,
prefix=add_prefix("attn", prefix),
workspace_buffer=workspace_buffer,
mapping=mapping,
mm_attention_backend=mm_attention_backend,
)
self.mlp = Qwen3VLVisionMLP(
dim,
intermediate_dim,
hidden_act=hidden_act,
bias=True,
quant_config=quant_config,
prefix=f"{prefix}.mlp",
mapping=mapping,
)
def forward(
self,
x: torch.Tensor,
cu_seqlens: torch.Tensor,
rotary_pos_emb_cos: torch.Tensor,
rotary_pos_emb_sin: torch.Tensor,
max_seqlen: int | None = None,
sequence_lengths: torch.Tensor | None = None,
) -> torch.Tensor:
hidden_states = self.norm1(x)
hidden_states = rearrange(hidden_states, "s b ... -> b s ...")
attn = self.attn(
hidden_states,
cu_seqlens=cu_seqlens,
rotary_pos_emb_cos=rotary_pos_emb_cos,
rotary_pos_emb_sin=rotary_pos_emb_sin,
max_seqlen=max_seqlen,
sequence_lengths=sequence_lengths,
)
attn = rearrange(attn, "b s ... -> s b ...")
x += attn
norm2 = self.norm2(x)
mlp = self.mlp(norm2)
x += mlp
return x
class Qwen3VLMoeVisionPatchMerger(nn.Module):
def __init__(
self,
dim: int,
context_dim: int,
padded_context_dim: int,
mapping: Mapping,
norm_layer: Callable[[int], nn.Module] | None = None,
spatial_merge_size: int = 2,
use_postshuffle_norm: bool = False,
quant_config: QuantizationConfig | None = None,
prefix: str = "",
) -> None:
super().__init__()
self.hidden_size = context_dim * (spatial_merge_size**2)
self.padded_context_dim = padded_context_dim * (spatial_merge_size**2)
self.use_postshuffle_norm = use_postshuffle_norm
if norm_layer is None:
norm_layer = partial(nn.LayerNorm, eps=1e-6)
self.norm = norm_layer(
self.hidden_size if use_postshuffle_norm else context_dim
)
vision = mapping.vision
self.linear_fc1 = ColumnParallelLinear(
self.hidden_size,
self.padded_context_dim,
bias=True,
quant_config=quant_config,
prefix=add_prefix("linear_fc1", prefix),
tp_size=vision.tp_size,
tp_rank=vision.tp_rank,
tp_group=vision.tp_group,
)
self.act_fn = nn.GELU()
self.linear_fc2 = RowParallelLinear(
self.padded_context_dim,
dim,
bias=True,
quant_config=quant_config,
prefix=add_prefix("linear_fc2", prefix),
tp_size=vision.tp_size,
tp_rank=vision.tp_rank,
tp_group=vision.tp_group,
reduce_results=True,
)
def forward(self, x: torch.Tensor) -> torch.Tensor:
if self.use_postshuffle_norm:
x = self.norm(x.view(-1, self.hidden_size))
else:
x = self.norm(x).view(-1, self.hidden_size)
x_parallel, _ = self.linear_fc1(x)
x_parallel = self.act_fn(x_parallel)
out, _ = self.linear_fc2(x_parallel)
return out
class Qwen3VLMoeVisionModel(nn.Module):
def __init__(
self,
vision_config: Qwen3VLVisionConfig,
mapping: Mapping,
norm_eps: float = 1e-6,
quant_config: QuantizationConfig | None = None,
prefix: str = "",
mm_attention_backend: str | None = None,
) -> None:
super().__init__()
vision = mapping.vision
self.mm_attention_backend = mm_attention_backend
self.hidden_size = vision_config.hidden_size
self.num_heads = vision_config.num_heads
self.num_position_embeddings = vision_config.num_position_embeddings
self.num_grid_per_side = int(self.num_position_embeddings**0.5)
self.spatial_merge_size = vision_config.spatial_merge_size
# layer indices whose outputs feed the deepstack mergers
self.deepstack_visual_indexes = vision_config.deepstack_visual_indexes
self.patch_embed = Qwen3VLVisionPatchEmbed(config=vision_config)
self.pos_embed = VocabParallelEmbedding(
self.num_position_embeddings,
self.hidden_size,
quant_config=quant_config,
tp_rank=vision.tp_rank,
tp_size=vision.tp_size,
tp_group=vision.tp_group,
prefix=add_prefix("pos_embed", prefix),
)
norm_layer = partial(nn.LayerNorm, eps=norm_eps)
head_dim = self.hidden_size // self.num_heads
self.rotary_pos_emb = get_rope(
head_size=head_dim,
rotary_dim=head_dim // 2,
max_position=8192,
base=10000.0,
is_neox_style=True,
)
workspace_buffer = None
if self.mm_attention_backend == "flashinfer_cudnn":
if torch.cuda.is_available():
ws_device = torch.device("cuda", torch.cuda.current_device())
else:
ws_device = self.device
workspace_buffer = torch.empty(
VIT_CUDNN_WORKSPACE_BYTES,
dtype=torch.uint8,
device=ws_device,
)
self.blocks = nn.ModuleList(
[
Qwen3VLVisionBlock(
dim=self.hidden_size,
num_heads=self.num_heads,
intermediate_dim=vision_config.intermediate_size,
head_size=head_dim,
hidden_act=vision_config.hidden_act,
norm_layer=norm_layer,
quant_config=quant_config,
prefix=add_prefix(f"blocks.{layer_idx}", prefix),
workspace_buffer=workspace_buffer,
mapping=mapping,
mm_attention_backend=mm_attention_backend,
)
for layer_idx in range(vision_config.depth)
]
)
self.merger = Qwen3VLMoeVisionPatchMerger(
dim=vision_config.out_hidden_size,
context_dim=self.hidden_size,
padded_context_dim=self.num_heads * head_dim,
norm_layer=norm_layer,
spatial_merge_size=self.spatial_merge_size,
quant_config=quant_config,
prefix=add_prefix("merger", prefix),
mapping=mapping,
)
self.deepstack_merger_list = nn.ModuleList(
[
Qwen3VLMoeVisionPatchMerger(
dim=vision_config.out_hidden_size,
context_dim=self.hidden_size,
padded_context_dim=self.num_heads * head_dim,
spatial_merge_size=self.spatial_merge_size,
use_postshuffle_norm=True,
norm_layer=norm_layer,
quant_config=quant_config,
prefix=add_prefix(f"deepstack_merger_list.{layer_idx}", prefix),
mapping=mapping,
)
for layer_idx in range(len(self.deepstack_visual_indexes))
]
)
self.tp_size = vision.tp_size
@property
def dtype(self) -> torch.dtype:
return self.patch_embed.proj.weight.dtype
@property
def device(self) -> torch.device:
return self.patch_embed.proj.weight.device
def rot_pos_emb(
self, grid_thw: list[list[int]]
) -> tuple[torch.Tensor, torch.Tensor]:
pos_ids = []
for t, h, w in grid_thw:
base = _rot_pos_ids(h, w, self.spatial_merge_size)
pos_ids.append(base if t == 1 else base.repeat(t, 1))
pos_ids = torch.cat(pos_ids, dim=0).to(self.device, non_blocking=True)
max_grid_size = max(max(h, w) for _, h, w in grid_thw)
cos, sin = self._get_rotary_cos_sin(max_grid_size)
cos_combined = cos[pos_ids].flatten(1)
sin_combined = sin[pos_ids].flatten(1)
return cos_combined, sin_combined
def _get_rotary_cos_sin(self, seqlen: int) -> tuple[torch.Tensor, torch.Tensor]:
cos_sin = self.rotary_pos_emb.cos_sin_cache[:seqlen].to(self.device)
return cos_sin.chunk(2, dim=-1)
def fast_pos_embed_interpolate_from_list(self, grid_thw):
num_grid_per_side = self.num_grid_per_side
m_size = self.spatial_merge_size
hidden_dim = self.pos_embed.embedding_dim
outputs = []
for t, h, w in grid_thw:
h_idxs = torch.linspace(
0, num_grid_per_side - 1, h, dtype=torch.float32, device=self.device
)
w_idxs = torch.linspace(
0, num_grid_per_side - 1, w, dtype=torch.float32, device=self.device
)
h_floor = h_idxs.to(torch.long)
w_floor = w_idxs.to(torch.long)
h_ceil = torch.clamp(h_floor + 1, max=num_grid_per_side - 1)
w_ceil = torch.clamp(w_floor + 1, max=num_grid_per_side - 1)
dh = h_idxs - h_floor
dw = w_idxs - w_floor
# Create meshgrid view for all h, w vars
dh_grid, dw_grid = torch.meshgrid(dh, dw, indexing="ij")
h_floor_grid, w_floor_grid = torch.meshgrid(h_floor, w_floor, indexing="ij")
h_ceil_grid, w_ceil_grid = torch.meshgrid(h_ceil, w_ceil, indexing="ij")
# original computation of weights
# w00 = (1 - dh_grid) * (1 - dw_grid)
# w01 = (1 - dh_grid) * dw_grid
# w10 = dh_grid * (1 - dw_grid)
# w11 = dh_grid * dw_grid
# we reuse w11 here to avoid duplicate
# dh_grid * dw_grid computation
w11 = dh_grid * dw_grid
w10 = dh_grid - w11
w01 = dw_grid - w11
w00 = 1 - dh_grid - w01
h_grid = torch.stack([h_floor_grid, h_floor_grid, h_ceil_grid, h_ceil_grid])
w_grid = torch.stack([w_floor_grid, w_ceil_grid, w_floor_grid, w_ceil_grid])
h_grid_idx = h_grid * num_grid_per_side
indices = (h_grid_idx + w_grid).reshape(4, -1)
weights = torch.stack([w00, w01, w10, w11], dim=0).reshape(4, -1, 1)
weights = weights.to(dtype=self.dtype)
embeds = self.pos_embed(indices)
embeds *= weights
combined = embeds.sum(dim=0)
combined = combined.reshape(
h // m_size, m_size, w // m_size, m_size, hidden_dim
)
combined = combined.permute(0, 2, 1, 3, 4).reshape(1, -1, hidden_dim)
repeated = combined.expand(t, -1, -1).reshape(-1, hidden_dim)
outputs.append(repeated)
return torch.cat(outputs, dim=0)
def compute_cudnn_batch_offsets_packed(
self,
token_cu_seqlens: np.ndarray,
*,
elem_per_token: int,
) -> np.ndarray:
"""
Build packed *element* indptrs for cuDNN prefill.
Input:
token_cu_seqlens: (B+1,) token indptr
elem_per_token: per-token element width on THIS TP rank
(usually hidden_size / attn_tp_size)
Output:
packed_offsets: (3 * (B_padded + 1),) int32
[qk_indptr, v_indptr, o_indptr] concatenated,
each indptr is (B_padded + 1,) in element units.
"""
if token_cu_seqlens.ndim != 1 or token_cu_seqlens.size < 2:
raise ValueError(
"token_cu_seqlens must be a 1D array with at least 2 entries."
)
B = int(token_cu_seqlens.size - 1)
B_padded = round_up_to_bucket(B, VIT_CUDNN_BATCH_BUCKETS)
# token indptr -> pad to (B_padded+1,) by appending total_tokens for extra empty sequences
token_indptr = token_cu_seqlens.astype(np.int64, copy=False) # (B+1,)
if B_padded != B:
pad = np.full((B_padded - B,), token_indptr[-1], dtype=token_indptr.dtype)
token_indptr = np.concatenate([token_indptr, pad], axis=0) # (B_padded+1,)
# convert token indptr -> element indptr
elem_indptr = (token_indptr * int(elem_per_token)).astype(
np.int32
) # (B_padded+1,)
# q/k/v/o in this vision encoder path share the same indptr
return np.concatenate([elem_indptr, elem_indptr, elem_indptr], axis=0)
def compute_cudnn_sequence_lengths_padded(
self,
token_cu_seqlens: np.ndarray,
) -> np.ndarray:
"""
token_cu_seqlens: (B+1,) token indptr
return: (B_padded,) token lengths (padded with 0)
"""
if token_cu_seqlens.ndim != 1 or token_cu_seqlens.size < 2:
raise ValueError(
"token_cu_seqlens must be a 1D array with at least 2 entries."
)
B = int(token_cu_seqlens.size - 1)
seq_lens = (token_cu_seqlens[1:] - token_cu_seqlens[:-1]).astype(
np.int32
) # (B,)
B_padded = round_up_to_bucket(B, VIT_CUDNN_BATCH_BUCKETS)
if B_padded != B:
pad = np.zeros((B_padded - B,), dtype=np.int32)
seq_lens = np.concatenate([seq_lens, pad], axis=0) # (B_padded,)
return seq_lens
def prepare_patch_embed(
self, x: torch.Tensor, grid_thw: torch.Tensor | list
) -> torch.Tensor:
"""Eager patch-embed (runs before the captured region): Conv patch embed
+ interpolated position embedding + the ``[s, 1, h]`` reshape the block
loop expects.
Kept eager (outside the capture-safe region) -- the interpolation does
host/numpy work.
"""
x = x.to(device=self.device, dtype=self.dtype)
x = self.patch_embed(x)
grid_thw_list = grid_thw if isinstance(grid_thw, list) else grid_thw.tolist()
x = x + self.fast_pos_embed_interpolate_from_list(grid_thw_list)
return x.unsqueeze(1)
def prepare_metadata(self, grid_thw: torch.Tensor | list) -> dict:
"""Eager metadata pass: rotary embeddings, cu_seqlens, sequence lengths,
and ``max_seqlen`` as a Python int.
Everything here involves a host sync or a data-dependent shape, so it
lives outside the capture-safe block loop. ``max_seqlen`` is
materialized as a plain int (CPU/numpy, no GPU sync) so the captured
block loop never hits the attention backend's ``.item()`` fallback.
"""
if isinstance(grid_thw, list):
grid_thw_list = grid_thw
grid_thw_np = np.array(grid_thw, dtype=np.int32)
else:
grid_thw_list = grid_thw.tolist()
grid_thw_np = grid_thw.cpu().numpy()
rotary_pos_emb_cos, rotary_pos_emb_sin = self.rot_pos_emb(grid_thw_list)
# ---- build token indptr (B+1,) ----
token_cu_seqlens = np.concatenate(
[
np.zeros(1, dtype=np.int32),
np.repeat(
grid_thw_np[:, 1] * grid_thw_np[:, 2], grid_thw_np[:, 0]
).cumsum(axis=0, dtype=np.int32),
]
)
real_seq_lens = token_cu_seqlens[1:] - token_cu_seqlens[:-1]
real_max_seqlen = int(real_seq_lens.max()) if real_seq_lens.size > 0 else 0
if self.mm_attention_backend == "flashinfer_cudnn":
# (B_padded,) token lengths
seq_lens_padded = self.compute_cudnn_sequence_lengths_padded(
token_cu_seqlens
)
# element-per-token width on this vision TP rank
elem_per_token = (
self.hidden_size // self.tp_size
) # == heads_per_rank * head_dim
# (3*(B_padded+1),) packed element indptrs
offsets_packed = self.compute_cudnn_batch_offsets_packed(
token_cu_seqlens,
elem_per_token=elem_per_token,
)
sequence_lengths = (
torch.from_numpy(seq_lens_padded)
.to(device=self.device, dtype=torch.int32, non_blocking=True)
.view(-1, 1, 1, 1)
) # match cuDNN test style
cu_seqlens = torch.from_numpy(offsets_packed).to(
device=self.device, dtype=torch.int32, non_blocking=True
)
max_seqlen = round_up_to_bucket(real_max_seqlen, VIT_CUDNN_SEQLEN_BUCKETS)
else:
sequence_lengths = None
cu_seqlens = torch.from_numpy(token_cu_seqlens).to(
device=self.device, dtype=torch.int32, non_blocking=True
)
max_seqlen = real_max_seqlen
return {
"cu_seqlens": cu_seqlens,
"rotary_pos_emb_cos": rotary_pos_emb_cos,
"rotary_pos_emb_sin": rotary_pos_emb_sin,
"max_seqlen": max_seqlen,
"sequence_lengths": sequence_lengths,
}
def forward_blocks(self, x: torch.Tensor, metadata: dict) -> torch.Tensor:
"""Capture-safe encoder body: block loop + deepstack mergers + merger.
No host syncs and no data-dependent control flow, so this region is
safe to record into a CUDA graph. ``metadata`` comes from
:meth:`prepare_metadata`; ``x`` from :meth:`prepare_patch_embed`.
"""
cu_seqlens = metadata["cu_seqlens"]
rotary_pos_emb_cos = metadata["rotary_pos_emb_cos"]
rotary_pos_emb_sin = metadata["rotary_pos_emb_sin"]
max_seqlen = metadata["max_seqlen"]
sequence_lengths = metadata["sequence_lengths"]
deepstack_feature_lists = []
num_deepstack_captured = 0
for layer_num, blk in enumerate(self.blocks):
x = blk(
x,
cu_seqlens=cu_seqlens,
rotary_pos_emb_cos=rotary_pos_emb_cos,
rotary_pos_emb_sin=rotary_pos_emb_sin,
max_seqlen=max_seqlen,
sequence_lengths=sequence_lengths,
)
if layer_num in self.deepstack_visual_indexes:
deepstack_feature = self.deepstack_merger_list[num_deepstack_captured](
x
)
deepstack_feature_lists.append(deepstack_feature)
num_deepstack_captured += 1
x = self.merger(x)
# [seq_len, out_hidden_size * (1 + depth_of_deepstack)]
return torch.cat([x] + deepstack_feature_lists, dim=1)