# Adopted from https://github.com/guandeh17/Self-Forcing # SPDX-License-Identifier: Apache-2.0 # # Copyright 2024-2025 The Alibaba Wan Team Authors. All rights reserved. from transformers.models.x_clip.modeling_x_clip import x_clip_loss from wan_5b.modules.attention import attention from wan_5b.modules.model import ( WanRMSNorm, rope_apply, WanLayerNorm, WanCrossAttention, rope_params, sinusoidal_embedding_1d, WanCrossAttention, flash_attention ) from torch.nn.attention.flex_attention import create_block_mask, flex_attention from diffusers.configuration_utils import ConfigMixin, register_to_config from torch.nn.attention.flex_attention import BlockMask from diffusers.models.modeling_utils import ModelMixin import os import torch.nn as nn import torch import math import torch.distributed as dist # wan 5b model compilation for flexattention flex_attention = torch.compile( flex_attention, dynamic=False, mode="max-autotune-no-cudagraphs") from utils.position_embedding_utils import ( compute_temporal_freqs as _compute_temporal_freqs, select_temporal_offset_for_sample, ) # iter-21: cache freqs_i across causal_rope_apply calls within a chunk. # All ~60 layer Q/K calls in one chunk share identical (f,h,w,start_frame, # t_scale,temporal_offset_i,method,original_seq_len) but recompute the same # concatenated freqs tensor each time. LRU keeps memory bounded. # NOTE: this cache holds tensors across torch.compile step boundaries which # is incompatible with cudagraphs (mode=reduce-overhead). If cudagraphs # path is enabled in the future, this cache must be removed alongside # refactoring of the KV cache scalar tensors (global_end_index, etc.). _FREQS_I_CACHE: "dict[tuple, torch.Tensor]" = {} _FREQS_I_CACHE_MAX = 16 # iter-21 + iter-41: cache is on by default (iter-21 win). Set # LLV2_FREQS_I_CACHE=0 to disable for future cudagraphs experiments (the # cache holds tensors created inside torch.compile that get marked as # cudagraph-pool memory; reading them on a later compile step crashes with # "accessing tensor output of CUDAGraphs that has been overwritten"). _FREQS_I_CACHE_ENABLED = os.environ.get("LLV2_FREQS_I_CACHE", "1") == "1" # iter-42: Triton fp32 RoPE kernel (utils/rope_triton.py). Default ON. # Replaces the fp64 complex view_as_complex × complex_freqs × view_as_real # chain with a single fused Triton kernel. Quality validated bit-exact at # bf16 (unit test agent/rope_unit_test.py: max|Δ|=7.8e-3 = single bf16 ULP). # Set LLV2_TRITON_ROPE=0 to revert to the fp64 path. # When enabled, _FREQS_I_CACHE stores (freqs_i_complex, cos_f32, sin_f32); # when disabled, stores (freqs_i_complex, None, None). _TRITON_ROPE_ENABLED = os.environ.get("LLV2_TRITON_ROPE", "1") == "1" # Cudagraph experiment only. Default OFF because the out-of-place temp-KV # construction removes mutated-input skips but is materially slower than the # in-place temporary cache update path. _CGRAPH_OUTPLACE_KV_ENABLED = os.environ.get("LLV2_CGRAPH_OUTPLACE_KV", "0") == "1" # iter-43/44: Triton fused adaLN-modulate kernel (utils/adaln_triton.py). # Default ON after iter-44 added `@triton.autotune` over (num_warps, num_stages). # iter-43 (no autotune) was FLAT vs iter-42 (median -1.0%, p90 +5.8%, total # identical) — fixed config beat the eager median but jitter on tail. # iter-44 (autotuned) is WIN: median -1.7%, p90 -1.6%, total -1.5%, FPS +1.5% # vs iter-42, quality in run-to-run noise floor (mean|Δ|=0.68 vs noise=0.69). # Unit test agent/adaln_unit_test.py: max|Δ|=3.1e-2 (1 bf16 ULP), mean=1.1e-3. # Set LLV2_TRITON_ADALN=0 to fall back to eager nn.LayerNorm + Python modulate. _TRITON_ADALN_ENABLED = os.environ.get("LLV2_TRITON_ADALN", "1") == "1" # iter-31: per-chunk Python-int metadata published by CausalWanModel.forward # so attention forwards can read Python ints without `.item()` graph breaks. # Single-thread inference assumption — overwritten before each model() call. _CURRENT_GRID_META: "dict[str, int]" = {} # iter-35: removed (LOST). Consolidating duplicate .item() reads caused # p90 latency to spike +10% — dynamo apparently traced more specialized # paths when local vars were used in branches vs fresh .item() reads each # time. Restored original .item() per-use pattern. def causal_rope_apply(x, grid_sizes, freqs, start_frame=0, t_scale=1.0, method="linear", original_seq_len=None, temporal_offset=0.0): n, c = x.size(2), x.size(3) // 2 # split freqs freqs = freqs.split([c - 2 * (c // 3), c // 3, c // 3], dim=1) # loop over samples output = [] # iter-47 (grad-safety fix): the Triton RoPE kernel (rope_apply_triton) is a # raw @triton.jit op with NO autograd backward — its output is graph-detached # (requires_grad=False). Running it under grad would silently sever gradients # to q/k (and their LoRA). Mirror the adaLN gate (see `use_triton_adaln`): # use Triton only when grad is OFF (inference / no_grad rollout steps); fall # back to the differentiable fp64-complex path whenever grad is ON (training). use_triton_rope = _TRITON_ROPE_ENABLED and not torch.is_grad_enabled() # iter-30: accept Python list/tuple to skip the .tolist() graph break. # Callers that already have Python ints (sink_grid, local_grid, window_grid_sizes) # now pass a plain list instead of `torch.tensor([[..]]).expand(...)`. if isinstance(grid_sizes, (list, tuple)): fwh_list = grid_sizes else: fwh_list = grid_sizes.tolist() for i, (f, h, w) in enumerate(fwh_list): seq_len = f * h * w # precompute multipliers — only needed for the fp64 complex path. # iter-42: skip the bf16→fp64 cast + view_as_complex when the Triton # kernel will be used (it consumes bf16 directly). # iter-47: gate on use_triton_rope (not the raw flag) so the complex x_i # IS precomputed whenever we fall back to the differentiable path (training). if not use_triton_rope: x_i = torch.view_as_complex(x[i, :seq_len].to(torch.float64).reshape( seq_len, n, -1, 2)) temporal_offset_i = select_temporal_offset_for_sample( temporal_offset, i, f, start_frame=start_frame) # iter-21: cache freqs_i. iter-41: gate the cache behind # LLV2_FREQS_I_CACHE=1 (default off). The cache stores tensors # created inside torch.compile, which cudagraph allocator considers # owned by the per-step memory pool — reading them on a later step # races with the pool's reuse. Disabling the cache unblocks # `mode=reduce-overhead` for cudagraphs; the recomputation cost is # tiny (60 layer calls × per-chunk concat ≈ 0.5% wall) compared to # the cudagraphs unlock potential. if _FREQS_I_CACHE_ENABLED: if torch.is_tensor(temporal_offset_i): if temporal_offset_i.ndim == 0: offset_key = float(temporal_offset_i.item()) else: offset_key = ("tensor", id(temporal_offset_i)) else: offset_key = float(temporal_offset_i) cache_key = ( f, h, w, start_frame, t_scale, method, original_seq_len, offset_key, x.device.type, x.device.index, use_triton_rope, # iter-47: separate cached repr for grad/no-grad ) cache_entry = _FREQS_I_CACHE.get(cache_key) else: cache_entry = None cache_key = None if cache_entry is None: temporal_freqs = _compute_temporal_freqs( freqs[0], f, start_frame, t_scale, x.device, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset_i) freqs_i_complex = torch.cat([ temporal_freqs.view(f, 1, 1, -1).expand(f, h, w, -1), freqs[1][:h].view(1, h, 1, -1).expand(f, h, w, -1), freqs[2][:w].view(1, 1, w, -1).expand(f, h, w, -1), ], dim=-1).reshape(seq_len, 1, -1) if use_triton_rope: # iter-42: store (cos, sin) fp32 derived once; freqs_i_complex # itself only kept for the legacy fp64 path. from utils.rope_triton import _split_complex_to_cos_sin cos_f32, sin_f32 = _split_complex_to_cos_sin(freqs_i_complex) cache_entry = (freqs_i_complex, cos_f32, sin_f32) else: cache_entry = (freqs_i_complex, None, None) if _FREQS_I_CACHE_ENABLED: if len(_FREQS_I_CACHE) >= _FREQS_I_CACHE_MAX: _FREQS_I_CACHE.pop(next(iter(_FREQS_I_CACHE))) _FREQS_I_CACHE[cache_key] = cache_entry freqs_i, cos_f32, sin_f32 = cache_entry # apply rotary embedding if use_triton_rope: # iter-42: Triton fp32 kernel — replaces the fp64 complex128 path. # iter-46: kernel takes full x[i] + seq_len and emits rotated-or- # passthrough output in a single launch, eliminating the # `.contiguous()` slice + outer `torch.cat`. Bit-exact preserved. from utils.rope_triton import rope_apply_triton x_i = rope_apply_triton(x[i], cos_f32, sin_f32, seq_len=seq_len) else: x_i = torch.view_as_real(x_i * freqs_i).flatten(2) x_i = torch.cat([x_i, x[i, seq_len:]]) # append to collection output.append(x_i) return torch.stack(output).type_as(x) class MultiShotT2VCrossAttention(WanCrossAttention): def forward(self, x, context, context_lens, is_teacher_forcing=False, crossattn_cache=None): r""" Args: x(Tensor): Shape [B, L1, C] context(Tensor): Shape [B * num_chunks, L2, C] context_lens(Tensor): Shape [B] or [B * num_chunks] crossattn_cache (List[dict], *optional*): Contains the cached key and value tensors for context embedding. """ # Original batch size (videos) b_orig, L1, C = x.size() n, d = self.num_heads, self.head_dim # Effective batch size for cross-attention (videos * chunks) b_ctx = context.size(0) assert b_ctx % b_orig == 0, f"context batch ({b_ctx}) must be a multiple of x batch ({b_orig})" num_chunks = b_ctx // b_orig # Prepare context_lens for [B * num_chunks] if needed if context_lens is not None and context_lens.numel() == b_orig: context_lens = context_lens.repeat_interleave(num_chunks) elif context_lens is not None: assert context_lens.numel() == b_ctx, \ f"context_lens must have length {b_orig} or {b_ctx}, got {context_lens.numel()}" # Helper to run standard cross-attention on a given x_chunk of shape [B * num_chunks, L_chunk, C] def _cross_attend(x_chunk): b_eff, L_chunk, _ = x_chunk.size() # compute query, key, value q = self.norm_q(self.q(x_chunk)).view(b_eff, -1, n, d) # iter-24: Bypass crossattn_cache. Cached K/V tensors escape the # cudagraph memory pool across torch.compile step boundaries and # block `mode=reduce-overhead`. Per-call recompute cost is tiny # (~1.7us / call in NVFP4 × ~11.5k calls/prompt ≈ 19 ms total), # for cudagraphs unlock of the 28% wall-time gap. k = self.norm_k(self.k(context)).view(b_eff, -1, n, d) v = self.v(context).view(b_eff, -1, n, d) # compute attention x_attn = flash_attention(q, k, v, k_lens=context_lens) # output projection x_attn = x_attn.flatten(2) x_attn = self.o(x_attn) return x_attn if not is_teacher_forcing: # ------------------------------- # Regular multi-shot: all tokens attend text, we just chunk along L1 # x: [B, L1, C] -> [B * num_chunks, L1 / num_chunks, C] # ------------------------------- assert L1 % num_chunks == 0, \ f"L1 ({L1}) must be divisible by num_chunks ({num_chunks})" tokens_per_chunk = L1 // num_chunks x_chunked = x.view(b_orig, num_chunks, tokens_per_chunk, C) x_chunked = x_chunked.reshape(b_ctx, tokens_per_chunk, C) x_attn = _cross_attend(x_chunked) # [B * num_chunks, tokens_per_chunk, C] # reshape back to [B, L1, C] x_attn = x_attn.view(b_orig, num_chunks, tokens_per_chunk, C) x_attn = x_attn.reshape(b_orig, L1, C) return x_attn # ------------------------------- # Teacher forcing: # x is typically [B, 2 * L_tf, C], where the first half is clean and # the second half is noisy. Apply multi-shot cross-attention to both # halves. # ------------------------------- assert L1 % 2 == 0, f"In teacher-forcing mode, L1 ({L1}) should be even." half = L1 // 2 x_clean = x[:, :half, :] # [B, L_tf, C] x_noisy = x[:, half:, :] # [B, L_tf, C] def _chunk_and_attend(x_part): L_part = x_part.size(1) assert L_part % num_chunks == 0, \ f"Segment length ({L_part}) must be divisible by num_chunks ({num_chunks})" tokens_per_chunk = L_part // num_chunks # [B, L_part, C] -> [B * num_chunks, L_part / num_chunks, C] x_chunked = x_part.view(b_orig, num_chunks, tokens_per_chunk, C) x_chunked = x_chunked.reshape(b_ctx, tokens_per_chunk, C) x_attn = _cross_attend(x_chunked) # [B * num_chunks, tokens_per_chunk, C] x_attn = x_attn.view(b_orig, num_chunks, tokens_per_chunk, C) x_attn = x_attn.reshape(b_orig, L_part, C) return x_attn x_clean_attn = _chunk_and_attend(x_clean) x_noisy_attn = _chunk_and_attend(x_noisy) # Reassemble the full sequence from cross-attended clean and noisy halves. x_out = torch.cat([x_clean_attn, x_noisy_attn], dim=1) # [B, L1, C] return x_out class CausalWanSelfAttention(nn.Module): def __init__(self, dim, num_heads, local_attn_size=-1, sink_size=0, qk_norm=True, eps=1e-6): assert dim % num_heads == 0 super().__init__() self.dim = dim self.num_heads = num_heads self.head_dim = dim // num_heads self.local_attn_size = local_attn_size if local_attn_size != -1 else 24 self.sink_size = sink_size self.global_sink_size = 0 self.qk_norm = qk_norm self.eps = eps self.max_attention_size = 24 * 880 if local_attn_size == -1 else local_attn_size * 880 # layers self.q = nn.Linear(dim, dim) self.k = nn.Linear(dim, dim) self.v = nn.Linear(dim, dim) self.o = nn.Linear(dim, dim) self.norm_q = WanRMSNorm(dim, eps=eps) if qk_norm else nn.Identity() self.norm_k = WanRMSNorm(dim, eps=eps) if qk_norm else nn.Identity() def forward( self, x, seq_lens, grid_sizes, freqs, block_mask, kv_cache=None, current_start=0, cache_start=None, t_scale=1.0, use_relative_rope=False, method="linear", original_seq_len=None, temporal_offset=0.0, ): r""" Args: x(Tensor): Shape [B, L, num_heads, C / num_heads] seq_lens(Tensor): Shape [B] grid_sizes(Tensor): Shape [B, 3], the second dimension contains (F, H, W) freqs(Tensor): Rope freqs, shape [1024, C / num_heads / 2] block_mask (BlockMask) t_scale (float): Temporal RoPE interpolation scale. <1.0 compresses positions. use_relative_rope (bool): If True, store raw K in cache and apply RoPE with window-relative positions at attention time. method (str): RoPE method. This release supports "linear". original_seq_len (int): Unused by the release linear RoPE path. temporal_offset (float): Multi-shot RoPE offset (shot_index * phi). """ b, s, n, d = *x.shape[:2], self.num_heads, self.head_dim if cache_start is None: cache_start = current_start # query, key, value function def qkv_fn(x): q = self.norm_q(self.q(x)).view(b, s, n, d) k = self.norm_k(self.k(x)).view(b, s, n, d) v = self.v(x).view(b, s, n, d) return q, k, v q, k, v = qkv_fn(x) if kv_cache is None: # Teacher-forcing training doubles sequence length with clean/noisy halves. is_tf = (s == seq_lens[0].item() * 2) if is_tf: q_chunk = torch.chunk(q, 2, dim=1) k_chunk = torch.chunk(k, 2, dim=1) roped_query = [] roped_key = [] # rope should be same for clean and noisy parts for ii in range(2): rq = rope_apply(q_chunk[ii], grid_sizes, freqs, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) rk = rope_apply(k_chunk[ii], grid_sizes, freqs, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) roped_query.append(rq) roped_key.append(rk) roped_query = torch.cat(roped_query, dim=1) roped_key = torch.cat(roped_key, dim=1) padded_length = math.ceil(q.shape[1] / 128) * 128 - q.shape[1] padded_roped_query = torch.cat( [roped_query, torch.zeros([q.shape[0], padded_length, q.shape[2], q.shape[3]], device=q.device, dtype=v.dtype)], dim=1 ) padded_roped_key = torch.cat( [roped_key, torch.zeros([k.shape[0], padded_length, k.shape[2], k.shape[3]], device=k.device, dtype=v.dtype)], dim=1 ) padded_v = torch.cat( [v, torch.zeros([v.shape[0], padded_length, v.shape[2], v.shape[3]], device=v.device, dtype=v.dtype)], dim=1 ) x = flex_attention( query=padded_roped_query.transpose(2, 1), key=padded_roped_key.transpose(2, 1), value=padded_v.transpose(2, 1), block_mask=block_mask ) x = x[:, :, :(-padded_length)] if padded_length > 0 else x x = x.transpose(2, 1) else: roped_query = rope_apply(q, grid_sizes, freqs, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) roped_key = rope_apply(k, grid_sizes, freqs, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) padded_length = math.ceil(q.shape[1] / 128) * 128 - q.shape[1] padded_roped_query = torch.cat( [roped_query, torch.zeros([q.shape[0], padded_length, q.shape[2], q.shape[3]], device=q.device, dtype=v.dtype)], dim=1 ) padded_roped_key = torch.cat( [roped_key, torch.zeros([k.shape[0], padded_length, k.shape[2], k.shape[3]], device=k.device, dtype=v.dtype)], dim=1 ) padded_v = torch.cat( [v, torch.zeros([v.shape[0], padded_length, v.shape[2], v.shape[3]], device=v.device, dtype=v.dtype)], dim=1 ) x = flex_attention( query=padded_roped_query.transpose(2, 1), key=padded_roped_key.transpose(2, 1), value=padded_v.transpose(2, 1), block_mask=block_mask ) x = x[:, :, :(-padded_length)] if padded_length > 0 else x x = x.transpose(2, 1) else: # iter-31: read Python ints from module-level dict (set by # CausalWanModel.forward) instead of `.item()` calls on # grid_sizes, removing 4 graph breaks per attention forward. if _CURRENT_GRID_META: frame_seqlen = _CURRENT_GRID_META["frame_seqlen"] num_new_frames = _CURRENT_GRID_META["num_new_frames"] h = _CURRENT_GRID_META["h"] w = _CURRENT_GRID_META["w"] else: frame_seqlen = math.prod(grid_sizes[0][1:]).item() num_new_frames = grid_sizes[0][0].item() h, w = grid_sizes[0][1].item(), grid_sizes[0][2].item() num_new_tokens = q.shape[1] current_end = current_start + num_new_tokens # iter-30: build Python-int grid once; pass to all rope_apply calls # below so they skip the .tolist() graph break. b = q.shape[0] grid_py = [(num_new_frames, h, w)] * b if not use_relative_rope: current_start_frame = current_start // frame_seqlen roped_query = causal_rope_apply( q, grid_py, freqs, start_frame=current_start_frame, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) roped_key = causal_rope_apply( k, grid_py, freqs, start_frame=current_start_frame, t_scale=t_scale, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset).type_as(v) key_to_cache = roped_key else: key_to_cache = k sink_tokens = self.sink_size * frame_seqlen global_sink_tokens = getattr(self, "global_sink_size", 0) * frame_seqlen is_quantized_cache = kv_cache.get("quantized", False) if is_quantized_cache: kv_cache_size = kv_cache["max_blocks"] * kv_cache["block_token_size"] else: kv_cache_size = kv_cache["k"].shape[1] # ----- global + multi-shot pinned-sink support ----- # Two protection mechanisms (independent, both optional): # * global_sink_tokens: first N frames are permanently anchored # (set via global_sink_size; never moves, always attended). # * pinned region (pinned_start/pinned_len): multi-shot sink put # on a scene cut. The pinned chunk lives at its original buffer # position; rolling shifts non-pinned data around it. # effective_sink = leading buffer prefix that rolling MUST keep: # pinned right after global (pinned_start == global_sink_tokens) # -> global_sink_tokens + pinned_len # pinned elsewhere (floating) # -> global_sink_tokens # no pinned # -> max(global_sink_tokens, sink_tokens) # legacy compat # iter-39: read pinned state from _CURRENT_GRID_META (published # once per chunk in CausalWanModel._forward_inference). Falls # back to `.item()` if the dict was not initialized (e.g. unit # test exercising the attention block directly). if _CURRENT_GRID_META and "pinned_start" in _CURRENT_GRID_META: pinned_start_val = _CURRENT_GRID_META["pinned_start"] pinned_len_val = _CURRENT_GRID_META["pinned_len"] else: pinned_start_t = kv_cache.get("pinned_start", None) pinned_len_val = 0 if pinned_start_t is not None and hasattr(pinned_start_t, 'item'): pinned_start_val = pinned_start_t.item() pinned_len_val = kv_cache["pinned_len"].item() else: pinned_start_val = -1 has_pinned = pinned_start_val >= 0 and pinned_len_val > 0 if has_pinned and pinned_start_val == global_sink_tokens: effective_sink = global_sink_tokens + pinned_len_val elif has_pinned: effective_sink = global_sink_tokens else: effective_sink = max(global_sink_tokens, sink_tokens) # iter-39: read cache indices from _CURRENT_GRID_META (published # by CausalWanModel._forward_inference) to avoid 6+ `.item()` # syncs per block forward. Falls back to .item() when the dict # is not initialized (direct attention-block unit tests). if _CURRENT_GRID_META and "global_end_index" in _CURRENT_GRID_META: _cache_global_end = _CURRENT_GRID_META["global_end_index"] _cache_local_end = _CURRENT_GRID_META["local_end_index"] else: _cache_global_end = kv_cache["global_end_index"].item() _cache_local_end = kv_cache["local_end_index"].item() cache_update_info = None if self.local_attn_size != -1 and (current_end > _cache_global_end) and ( num_new_tokens + _cache_local_end > kv_cache_size): num_evicted_tokens = num_new_tokens + _cache_local_end - kv_cache_size num_rolled_tokens = _cache_local_end - num_evicted_tokens - effective_sink local_end_index = _cache_local_end + current_end - \ _cache_global_end - num_evicted_tokens local_start_index = local_end_index - num_new_tokens if is_quantized_cache: from utils.quant import dequantize_kv_cache, k_smooth max_blks = int(kv_cache["max_blocks"]) blk_sz = int(kv_cache["block_token_size"]) cache_k = dequantize_kv_cache( kv_cache["k"], max_blks, self.num_heads, blk_sz, v.dtype, v.device ) cache_v = dequantize_kv_cache( kv_cache["v"], max_blks, self.num_heads, blk_sz, v.dtype, v.device ) new_k_for_cache = k_smooth(key_to_cache) else: cache_k = kv_cache["k"] cache_v = kv_cache["v"] new_k_for_cache = key_to_cache if _CGRAPH_OUTPLACE_KV_ENABLED: # Cudagraph experiment: build the post-roll cache view # out-of-place. Slice assignment here forces Inductor # cudagraph partitions to mutate inputs. temp_k = torch.cat([ cache_k[:, :effective_sink], cache_k[:, effective_sink + num_evicted_tokens: effective_sink + num_evicted_tokens + num_rolled_tokens], new_k_for_cache, ], dim=1) temp_v = torch.cat([ cache_v[:, :effective_sink], cache_v[:, effective_sink + num_evicted_tokens: effective_sink + num_evicted_tokens + num_rolled_tokens], v, ], dim=1) else: temp_k = cache_k if is_quantized_cache else cache_k.clone() temp_v = cache_v if is_quantized_cache else cache_v.clone() temp_k[:, effective_sink:effective_sink + num_rolled_tokens] = \ temp_k[:, effective_sink + num_evicted_tokens:effective_sink + num_evicted_tokens + num_rolled_tokens].clone() temp_v[:, effective_sink:effective_sink + num_rolled_tokens] = \ temp_v[:, effective_sink + num_evicted_tokens:effective_sink + num_evicted_tokens + num_rolled_tokens].clone() temp_k[:, local_start_index:local_end_index] = new_k_for_cache temp_v[:, local_start_index:local_end_index] = v # When pinned is "floating" (lives outside effective_sink), the # rolling shifted non-pinned data left by num_evicted_tokens; # the pinned anchor must follow that shift to keep tracking the # same data. When pinned sits inside effective_sink (i.e. right # after the global region), it is part of the protected prefix # and rolling does not move it. pinned_shift = num_evicted_tokens if (has_pinned and pinned_start_val >= effective_sink) else 0 cache_update_info = { "action": "roll_and_insert", "sink_tokens": effective_sink, "num_rolled_tokens": num_rolled_tokens, "num_evicted_tokens": num_evicted_tokens, "local_start_index": local_start_index, "local_end_index": local_end_index, "new_k": key_to_cache, "new_v": v, "current_end": current_end, "pinned_shift": pinned_shift, } else: # iter-39: reuse the dict-cached scalars from above. local_end_index = _cache_local_end + current_end - _cache_global_end local_start_index = local_end_index - num_new_tokens if is_quantized_cache: from utils.quant import dequantize_kv_cache, k_smooth new_k_for_cache = k_smooth(key_to_cache) if local_start_index == 0: temp_k = new_k_for_cache temp_v = v else: max_blks = int(kv_cache["max_blocks"]) blk_sz = int(kv_cache["block_token_size"]) cache_k = dequantize_kv_cache( kv_cache["k"], max_blks, self.num_heads, blk_sz, v.dtype, v.device ) cache_v = dequantize_kv_cache( kv_cache["v"], max_blks, self.num_heads, blk_sz, v.dtype, v.device ) if _CGRAPH_OUTPLACE_KV_ENABLED: temp_k = torch.cat([cache_k[:, :local_start_index], new_k_for_cache], dim=1) temp_v = torch.cat([cache_v[:, :local_start_index], v], dim=1) else: temp_k = cache_k temp_v = cache_v if not _CGRAPH_OUTPLACE_KV_ENABLED: temp_k[:, local_start_index:local_end_index] = new_k_for_cache temp_v[:, local_start_index:local_end_index] = v else: if _CGRAPH_OUTPLACE_KV_ENABLED: temp_k = torch.cat([kv_cache["k"][:, :local_start_index], key_to_cache], dim=1) temp_v = torch.cat([kv_cache["v"][:, :local_start_index], v], dim=1) else: temp_k = kv_cache["k"].clone() temp_v = kv_cache["v"].clone() temp_k[:, local_start_index:local_end_index] = key_to_cache temp_v[:, local_start_index:local_end_index] = v cache_update_info = { "action": "direct_insert", "local_start_index": local_start_index, "local_end_index": local_end_index, "new_k": key_to_cache, "new_v": v, "current_end": current_end, "pinned_shift": 0, } window_start = max(0, local_end_index - self.max_attention_size) # Build the K/V actually attended over. # Cases: # (a) prepend_sink : effective_sink > 0 and out of window # -> prepend [:effective_sink] (covers global # and any pinned-merged-to-front) # (b) prepend_pinned: a floating pinned region (pinned_start # >= effective_sink) lives outside the window # -> additionally prepend that pinned slice # (c) otherwise : plain sliding window # Note (a) and (b) are not mutually exclusive: when global is # enabled AND there is a separate floating pinned region outside # the window, both prefixes must be prepended. prepend_sink = effective_sink > 0 and window_start > 0 prepend_pinned = ( has_pinned and pinned_start_val >= effective_sink and pinned_start_val < window_start ) if prepend_sink and prepend_pinned: # [global+sink] + [pinned] + [local window] extra = effective_sink + pinned_len_val effective_local_size = self.max_attention_size - extra local_window_start = max(effective_sink, local_end_index - effective_local_size) window_k = torch.cat([ temp_k[:, :effective_sink], temp_k[:, pinned_start_val:pinned_start_val + pinned_len_val], temp_k[:, local_window_start:local_end_index], ], dim=1) window_v = torch.cat([ temp_v[:, :effective_sink], temp_v[:, pinned_start_val:pinned_start_val + pinned_len_val], temp_v[:, local_window_start:local_end_index], ], dim=1) elif prepend_sink: effective_local_size = self.max_attention_size - effective_sink local_window_start = max(effective_sink, local_end_index - effective_local_size) window_k = torch.cat([temp_k[:, :effective_sink], temp_k[:, local_window_start:local_end_index]], dim=1) window_v = torch.cat([temp_v[:, :effective_sink], temp_v[:, local_window_start:local_end_index]], dim=1) elif prepend_pinned: effective_local_size = self.max_attention_size - pinned_len_val local_window_start = max(0, local_end_index - effective_local_size) window_k = torch.cat( [temp_k[:, pinned_start_val:pinned_start_val + pinned_len_val], temp_k[:, local_window_start:local_end_index]], dim=1) window_v = torch.cat( [temp_v[:, pinned_start_val:pinned_start_val + pinned_len_val], temp_v[:, local_window_start:local_end_index]], dim=1) else: window_k = temp_k[:, window_start:local_end_index] window_v = temp_v[:, window_start:local_end_index] if use_relative_rope: if prepend_sink: # Sink and local window tokens get separate RoPE in a # virtual contiguous layout: [sink_frames | local_frames]. sink_frame_count = effective_sink // frame_seqlen local_tokens = window_k.shape[1] - effective_sink local_frame_count = local_tokens // frame_seqlen combined_frames = sink_frame_count + local_frame_count # iter-30: pass Python list instead of expanded tensor; # causal_rope_apply skips .tolist() graph break this way. sink_grid = [(sink_frame_count, h, w)] * b roped_sink_k = causal_rope_apply( window_k[:, :effective_sink], sink_grid, freqs, start_frame=0, t_scale=t_scale, method=method, original_seq_len=original_seq_len, ).type_as(v) local_grid = [(local_frame_count, h, w)] * b roped_local_k = causal_rope_apply( window_k[:, effective_sink:], local_grid, freqs, start_frame=sink_frame_count, t_scale=t_scale, method=method, original_seq_len=original_seq_len, ).type_as(v) roped_window_k = torch.cat([roped_sink_k, roped_local_k], dim=1) q_start_frame = combined_frames - num_new_frames roped_query = causal_rope_apply( q, grid_py, freqs, start_frame=q_start_frame, t_scale=t_scale, method=method, original_seq_len=original_seq_len, ).type_as(v) else: window_tokens = window_k.shape[1] window_frames = window_tokens // frame_seqlen # iter-30: Python list to skip .tolist() break. window_grid_sizes = [(window_frames, h, w)] * b roped_window_k = causal_rope_apply( window_k, window_grid_sizes, freqs, start_frame=0, t_scale=t_scale, method=method, original_seq_len=original_seq_len, ).type_as(v) q_start_frame = window_frames - num_new_frames roped_query = causal_rope_apply( q, grid_py, freqs, start_frame=q_start_frame, t_scale=t_scale, method=method, original_seq_len=original_seq_len, ).type_as(v) x = attention(roped_query, roped_window_k, window_v) else: x = attention(roped_query, window_k, window_v) # output x = x.flatten(2) x = self.o(x) # Return both output and cache update info if kv_cache is not None: return x, (current_end, local_end_index, cache_update_info) else: return x class CausalWanAttentionBlock(nn.Module): def __init__(self, dim, ffn_dim, num_heads, local_attn_size=-1, sink_size=0, qk_norm=True, cross_attn_norm=False, eps=1e-6): super().__init__() self.dim = dim self.ffn_dim = ffn_dim self.num_heads = num_heads self.local_attn_size = local_attn_size self.qk_norm = qk_norm self.cross_attn_norm = cross_attn_norm self.eps = eps # layers self.norm1 = WanLayerNorm(dim, eps) self.self_attn = CausalWanSelfAttention(dim, num_heads, local_attn_size, sink_size, qk_norm, eps) self.norm3 = WanLayerNorm( dim, eps, elementwise_affine=True) if cross_attn_norm else nn.Identity() self.cross_attn = MultiShotT2VCrossAttention(dim, num_heads, (-1, -1), qk_norm, eps) self.norm2 = WanLayerNorm(dim, eps) self.ffn = nn.Sequential( nn.Linear(dim, ffn_dim), nn.GELU(approximate='tanh'), nn.Linear(ffn_dim, dim)) # modulation self.modulation = nn.Parameter(torch.randn(1, 6, dim) / dim**0.5) def forward( self, x, e, seq_lens, grid_sizes, freqs, context, context_lens, block_mask, kv_cache=None, crossattn_cache=None, current_start=0, cache_start=None, t_scale=1.0, use_relative_rope=False, method="linear", original_seq_len=None, temporal_offset=0.0, ): r""" Args: x(Tensor): Shape [B, L, C] e(Tensor): Shape [B, F, 6, C] seq_lens(Tensor): Shape [B], length of each sequence in batch grid_sizes(Tensor): Shape [B, 3], the second dimension contains (F, H, W) freqs(Tensor): Rope freqs, shape [1024, C / num_heads / 2] t_scale (float): Temporal RoPE interpolation scale. <1.0 compresses positions. use_relative_rope (bool): If True, use window-relative RoPE positions in KV cache path. method (str): RoPE method. This release supports "linear". original_seq_len (int): Unused by the release linear RoPE path. temporal_offset (float): Multi-shot RoPE offset (shot_index * phi). """ num_frames, frame_seqlen = e.shape[1], x.shape[1] // e.shape[1] e = (self.modulation.unsqueeze(1) + e).chunk(6, dim=2) use_triton_adaln = _TRITON_ADALN_ENABLED and not torch.is_grad_enabled() # self-attention if use_triton_adaln: # iter-43: fused LayerNorm + (1+e[1])*x + e[0] in one Triton kernel. from utils.adaln_triton import adaln_modulate_triton modulated_x = adaln_modulate_triton( x, e[1], e[0], frame_seqlen, eps=self.norm1.eps, add_one_to_scale=True, ) else: modulated_x = (self.norm1(x).unflatten(dim=1, sizes=(num_frames, frame_seqlen)) * (1 + e[1]) + e[0]).flatten(1, 2) self_attn_result = self.self_attn( modulated_x, seq_lens, grid_sizes, freqs, block_mask, kv_cache, current_start, cache_start, t_scale=t_scale, use_relative_rope=use_relative_rope, method=method, original_seq_len=original_seq_len, temporal_offset=temporal_offset) if kv_cache is not None: y, cache_update_info = self_attn_result else: y = self_attn_result cache_update_info = None x = x + (y.unflatten(dim=1, sizes=(num_frames, frame_seqlen)) * e[2]).flatten(1, 2) # cross-attention & ffn function # iter-40: avoid `seq_lens[0].item()` graph break. seq_lens[0] equals # num_new_frames * frame_seqlen at inference time, both of which are # Python ints already in _CURRENT_GRID_META (published by iter-31). # `is_tf` is True only in teacher-forcing training where x.shape[1] # is the doubled (clean+noisy) sequence — never at inference. if _CURRENT_GRID_META and "frame_seqlen" in _CURRENT_GRID_META: seq_len_py = ( _CURRENT_GRID_META["frame_seqlen"] * _CURRENT_GRID_META["num_new_frames"] ) is_tf = (x.shape[1] == seq_len_py * 2) else: is_tf = (x.shape[1] == seq_lens[0].item() * 2) def cross_attn_ffn(x, context, context_lens, e, crossattn_cache=None): x = x + self.cross_attn(self.norm3(x), context, context_lens, crossattn_cache=crossattn_cache, is_teacher_forcing=is_tf) if use_triton_adaln: # iter-43: fused LayerNorm + (1+e[4])*x + e[3] in one Triton kernel. from utils.adaln_triton import adaln_modulate_triton ffn_in = adaln_modulate_triton( x, e[4], e[3], frame_seqlen, eps=self.norm2.eps, add_one_to_scale=True, ) else: ffn_in = (self.norm2(x).unflatten(dim=1, sizes=(num_frames, frame_seqlen)) * (1 + e[4]) + e[3]).flatten(1, 2) y = self.ffn(ffn_in) x = x + (y.unflatten(dim=1, sizes=(num_frames, frame_seqlen)) * e[5]).flatten(1, 2) return x x = cross_attn_ffn(x, context, context_lens, e, crossattn_cache) if cache_update_info is not None: # cache_update_info is already formatted as # (current_end, local_end_index, cache_update_info). return x, cache_update_info else: return x class CausalHead(nn.Module): def __init__(self, dim, out_dim, patch_size, eps=1e-6): super().__init__() self.dim = dim self.out_dim = out_dim self.patch_size = patch_size self.eps = eps # layers out_dim = math.prod(patch_size) * out_dim self.norm = WanLayerNorm(dim, eps) self.head = nn.Linear(dim, out_dim) # modulation self.modulation = nn.Parameter(torch.randn(1, 2, dim) / dim**0.5) def forward(self, x, e): r""" Args: x(Tensor): Shape [B, L1, C] e(Tensor): Shape [B, F, 1, C] """ num_frames, frame_seqlen = e.shape[1], x.shape[1] // e.shape[1] e = (self.modulation.unsqueeze(1) + e).chunk(2, dim=2) x = (self.head(self.norm(x).unflatten(dim=1, sizes=(num_frames, frame_seqlen)) * (1 + e[1]) + e[0])) return x class CausalWanModel(ModelMixin, ConfigMixin): r""" Wan diffusion backbone supporting both text-to-video and image-to-video with causal attention. """ ignore_for_config = [ 'patch_size', 'cross_attn_norm', 'qk_norm', 'text_dim', 'window_size' ] _no_split_modules = ['WanAttentionBlock'] _supports_gradient_checkpointing = True @register_to_config def __init__(self, model_type='t2v', patch_size=(1, 2, 2), text_len=512, in_dim=16, dim=2048, ffn_dim=8192, freq_dim=256, text_dim=4096, out_dim=16, num_heads=16, num_layers=32, local_attn_size=-1, sink_size=0, num_frame_per_block=1, qk_norm=True, cross_attn_norm=True, eps=1e-6): r""" Initialize the diffusion model backbone. Args: model_type (`str`, *optional*, defaults to 't2v'): Model variant - 't2v' (text-to-video), 'i2v' (image-to-video), or 'ti2v' patch_size (`tuple`, *optional*, defaults to (1, 2, 2)): 3D patch dimensions for video embedding (t_patch, h_patch, w_patch) text_len (`int`, *optional*, defaults to 512): Fixed length for text embeddings in_dim (`int`, *optional*, defaults to 16): Input video channels (C_in) dim (`int`, *optional*, defaults to 2048): Hidden dimension of the transformer ffn_dim (`int`, *optional*, defaults to 8192): Intermediate dimension in feed-forward network freq_dim (`int`, *optional*, defaults to 256): Dimension for sinusoidal time embeddings text_dim (`int`, *optional*, defaults to 4096): Input dimension for text embeddings out_dim (`int`, *optional*, defaults to 16): Output video channels (C_out) num_heads (`int`, *optional*, defaults to 16): Number of attention heads num_layers (`int`, *optional*, defaults to 32): Number of transformer blocks local_attn_size (`int`, *optional*, defaults to -1): Window size for temporal local attention (-1 indicates global attention) sink_size (`int`, *optional*, defaults to 0): Size of the attention sink, we keep the first `sink_size` frames unchanged when rolling the KV cache qk_norm (`bool`, *optional*, defaults to True): Enable query/key normalization cross_attn_norm (`bool`, *optional*, defaults to False): Enable cross-attention normalization eps (`float`, *optional*, defaults to 1e-6): Epsilon value for normalization layers """ super().__init__() assert model_type in ['t2v', 'i2v', 'ti2v'] self.model_type = model_type self.patch_size = patch_size self.text_len = text_len self.in_dim = in_dim self.dim = dim self.ffn_dim = ffn_dim self.freq_dim = freq_dim self.text_dim = text_dim self.out_dim = out_dim self.num_heads = num_heads self.num_layers = num_layers self.local_attn_size = local_attn_size self.sink_size = sink_size self.qk_norm = qk_norm self.cross_attn_norm = cross_attn_norm self.eps = eps # embeddings self.patch_embedding = nn.Conv3d( in_dim, dim, kernel_size=patch_size, stride=patch_size) self.text_embedding = nn.Sequential( nn.Linear(text_dim, dim), nn.GELU(approximate='tanh'), nn.Linear(dim, dim)) self.time_embedding = nn.Sequential( nn.Linear(freq_dim, dim), nn.SiLU(), nn.Linear(dim, dim)) self.time_projection = nn.Sequential( nn.SiLU(), nn.Linear(dim, dim * 6)) # blocks self.blocks = nn.ModuleList([ CausalWanAttentionBlock(dim, ffn_dim, num_heads, local_attn_size, sink_size, qk_norm, cross_attn_norm, eps) for _ in range(num_layers) ]) # head self.head = CausalHead(dim, out_dim, patch_size, eps) # buffers (don't use register_buffer otherwise dtype will be changed in to()) assert (dim % num_heads) == 0 and (dim // num_heads) % 2 == 0 d = dim // num_heads self.freqs = torch.cat([ rope_params(1024, d - 4 * (d // 6)), rope_params(1024, 2 * (d // 6)), rope_params(1024, 2 * (d // 6)) ], dim=1) # initialize weights self.init_weights() self.gradient_checkpointing = False self.block_mask = None self._block_mask_batch_size = 0 self.num_frame_per_block = num_frame_per_block self.independent_first_frame = False self.t_scale = 1.0 self.use_relative_rope = False self.rope_method = "linear" self.original_seq_len = None self.rope_temporal_offset = 0.0 self.kv_quant_config = None def _set_gradient_checkpointing(self, module, value=False): self.gradient_checkpointing = value @staticmethod def _prepare_blockwise_causal_attn_mask_i2v( device: torch.device | str, num_frames: int = 31, frame_seqlen: int = 880, num_frame_per_block=3, batch_size=None, ) -> BlockMask: """ we will divide the token sequence into the following format [1 latent frame] [N latent frame] ... [N latent frame] The first frame is separated out to support I2V generation We use flexattention to construct the attention mask """ total_length = num_frames * frame_seqlen # we do right padding to get to a multiple of 128 padded_length = math.ceil(total_length / 128) * 128 - total_length ends = torch.zeros(total_length + padded_length, device=device, dtype=torch.long) # special handling for the first frame ends[:frame_seqlen] = frame_seqlen # Block-wise causal mask will attend to all elements that are before the end of the current chunk frame_indices = torch.arange( start=frame_seqlen, end=total_length, step=frame_seqlen * num_frame_per_block, device=device ) for idx, tmp in enumerate(frame_indices): ends[tmp:tmp + frame_seqlen * num_frame_per_block] = tmp + \ frame_seqlen * num_frame_per_block def attention_mask(b, h, q_idx, kv_idx): is_real_q = q_idx < total_length is_real_k = kv_idx < total_length return (q_idx == kv_idx) | (is_real_q & is_real_k & (kv_idx < ends[q_idx])) block_mask = create_block_mask(attention_mask, B=batch_size, H=None, Q_LEN=total_length + padded_length, KV_LEN=total_length + padded_length, _compile=False, device=device) return block_mask @staticmethod def _prepare_blockwise_causal_attn_mask( device: torch.device | str, num_frames: int = 31, frame_seqlen: int = 880, num_frame_per_block=1, batch_size=None, ) -> BlockMask: """ Block-wise causal mask. The mask is defined only by the AR chunk size: a token can attend to all tokens before the end of its current num_frame_per_block chunk. """ print(f"num_frame_per_block: {num_frame_per_block}") total_length = num_frames * frame_seqlen padded_length = math.ceil(total_length / 128) * 128 - total_length block_size = frame_seqlen * num_frame_per_block def attention_mask(b, h, q_idx, kv_idx): # Apply only to real tokens in [0, total_length); padding keeps only # self-loops. is_real_q = q_idx < total_length is_real_k = kv_idx < total_length # End position of the block containing the current token. current_block_end = ((q_idx // block_size) + 1) * block_size clean_mask = is_real_q & is_real_k & (kv_idx < current_block_end) eye_mask = q_idx == kv_idx return eye_mask | clean_mask block_mask = create_block_mask( attention_mask, B=batch_size, H=None, Q_LEN=total_length + padded_length, KV_LEN=total_length + padded_length, _compile=True, device=device, ) import torch.distributed as dist if not dist.is_initialized() or dist.get_rank() == 0: print( f" cache a block wise causal mask with block size of {num_frame_per_block} frames" ) print(block_mask) return block_mask @staticmethod def _prepare_teacher_forcing_mask( device: torch.device | str, num_frames: int = 31, frame_seqlen: int = 880, num_frame_per_block: int = 1, batch_size: int | None = None, ): total_length = num_frames * frame_seqlen * 2 padded_length = math.ceil(total_length / 128) * 128 - total_length clean_ends = num_frames * frame_seqlen attention_block_size = frame_seqlen * num_frame_per_block # Use pure mathematical coordinates; do not introduce external tensor # lookup tables here. def attention_mask(b, h, q_idx, kv_idx): is_real_q = q_idx < total_length is_real_k = kv_idx < total_length # ========================================== # 1. Clean-frame mask. # ========================================== is_clean_q = q_idx < clean_ends # End position of the block containing the current token. clean_block_idx = q_idx // attention_block_size current_clean_block_end = (clean_block_idx + 1) * attention_block_size clean_mask = ( is_clean_q & (kv_idx < current_clean_block_end) ) # ========================================== # 2. Noisy-frame mask. # ========================================== is_noisy_q = q_idx >= clean_ends noisy_rel_idx = q_idx - clean_ends block_index = noisy_rel_idx // attention_block_size # C1: noisy tokens in the same block. noisy_block_start = clean_ends + (block_index * attention_block_size) noisy_block_end = noisy_block_start + attention_block_size C1 = (kv_idx >= noisy_block_start) & (kv_idx < noisy_block_end) # C2: clean context tokens from previous blocks. context_end_for_noisy = block_index * attention_block_size C2 = kv_idx < context_end_for_noisy noise_mask = is_noisy_q & (C1 | C2) # ========================================== # 3. Final merge. # ========================================== eye_mask = q_idx == kv_idx return eye_mask | (is_real_q & is_real_k & (clean_mask | noise_mask)) # _compile=True is required here. Triton compiles the mathematical # formula above directly into a memory-efficient block-sparse matrix. block_mask = create_block_mask( attention_mask, B=batch_size, H=None, Q_LEN=total_length + padded_length, KV_LEN=total_length + padded_length, _compile=True, device=device, ) import torch.distributed as dist if not dist.is_initialized() or dist.get_rank() == 0: print(block_mask) return block_mask @staticmethod def _prepare_teacher_forcing_mask_natural( device: torch.device | str, num_frames: int, frame_seqlen: int, num_frame_per_block: int = 1, sp_size: int = 1, batch_size: int | None = None, ): """Teacher-Forcing attention mask for the *natural* interleaved layout produced directly by `all_to_all(scatter=head, gather=seq)`: [r0_clean, r0_noisy, r1_clean, r1_noisy, ..., r_{sp-1}_clean, r_{sp-1}_noisy] Semantically equivalent to :func:`_prepare_teacher_forcing_mask` (which assumes the [all_clean; all_noisy] layout), but the mask decodes ``is_noisy`` and ``global_frame`` directly from the token index, so :func:`distributed_flex_attention` no longer has to reshape/permute tokens after all_to_all. """ assert num_frames % sp_size == 0, ( f"num_frames ({num_frames}) must be divisible by sp_size ({sp_size}) " f"for natural TF layout" ) F_local = num_frames // sp_size clean_half = F_local * frame_seqlen # per-rank, clean side per_rank_len = 2 * clean_half # per-rank, clean + noisy total_length = num_frames * frame_seqlen * 2 padded_length = math.ceil(total_length / 128) * 128 - total_length def attention_mask(b, h, q_idx, kv_idx): is_real_q = q_idx < total_length is_real_k = kv_idx < total_length # ---- decode q ---- r_q = q_idx // per_rank_len in_rank_q = q_idx % per_rank_len is_noisy_q = in_rank_q >= clean_half side_q = in_rank_q % clean_half # offset within clean/noisy half global_f_q = r_q * F_local + side_q // frame_seqlen block_q = global_f_q // num_frame_per_block # ---- decode k ---- r_k = kv_idx // per_rank_len in_rank_k = kv_idx % per_rank_len is_noisy_k = in_rank_k >= clean_half side_k = in_rank_k % clean_half global_f_k = r_k * F_local + side_k // frame_seqlen block_k = global_f_k // num_frame_per_block # 1. clean_q -> clean_k: blockwise causal. clean2clean = ( (~is_noisy_q) & (~is_noisy_k) & (block_k <= block_q) ) # 2. noisy_q -> clean_k: strictly earlier blocks. noisy2clean = ( is_noisy_q & (~is_noisy_k) & (block_k < block_q) ) # 3. noisy_q -> noisy_k: only tokens within the same block. noisy2noisy = ( is_noisy_q & is_noisy_k & (block_k == block_q) ) eye_mask = q_idx == kv_idx return eye_mask | ( is_real_q & is_real_k & (clean2clean | noisy2clean | noisy2noisy) ) block_mask = create_block_mask( attention_mask, B=batch_size, H=None, Q_LEN=total_length + padded_length, KV_LEN=total_length + padded_length, _compile=True, device=device, ) import torch.distributed as dist if not dist.is_initialized() or dist.get_rank() == 0: print( f"[TF mask natural] sp_size={sp_size} F_local={F_local} " f"clean_half={clean_half} per_rank_len={per_rank_len} " f"total_length={total_length} " f"num_frame_per_block={num_frame_per_block}" ) print(block_mask) return block_mask def _apply_cache_updates(self, kv_cache, cache_update_infos): """ Applies cache updates collected from multiple blocks. Args: kv_cache: List of cache dictionaries for each block cache_update_infos: List of (block_index, cache_update_info) tuples """ for block_index, (current_end, local_end_index, update_info) in cache_update_infos: if update_info is not None: cache = kv_cache[block_index] is_quantized = cache.get("quantized", False) if update_info["action"] == "roll_and_insert": # Apply the rolling update. sink_tokens = update_info["sink_tokens"] num_rolled_tokens = update_info["num_rolled_tokens"] num_evicted_tokens = update_info["num_evicted_tokens"] local_start_index = update_info["local_start_index"] local_end_index = update_info["local_end_index"] new_k = update_info["new_k"] new_v = update_info["new_v"] if is_quantized: from utils.quant import copy_quantized_into, quantize_to_fp4 blk_sz = int(cache["block_token_size"]) sink_blks = sink_tokens // blk_sz evict_blks = num_evicted_tokens // blk_sz roll_blks = num_rolled_tokens // blk_sz # iter-26: in-place copy into pre-allocated cache # slots instead of replacing the QuantizedTensor # reference. Required to unblock cudagraphs (the # fresh QT returned by quantize_to_fp4 lives in # the cudagraph memory pool; copying its data into # the persistent slot buffer breaks that escape). for i in range(roll_blks): src = sink_blks + evict_blks + i dst = sink_blks + i copy_quantized_into(cache["k"][dst], cache["k"][src]) copy_quantized_into(cache["v"][dst], cache["v"][src]) start_blk = local_start_index // blk_sz n_insert_blks = (local_end_index - local_start_index) // blk_sz head_dim = new_k.shape[-1] for bi in range(n_insert_blks): blk_idx = start_blk + bi ts = bi * blk_sz te = ts + blk_sz k_block = new_k[0, ts:te, :, :].reshape(-1, head_dim).contiguous() v_block = new_v[0, ts:te, :, :].reshape(-1, head_dim).contiguous() copy_quantized_into( cache["k"][blk_idx], quantize_to_fp4(k_block, self.kv_quant_config), ) copy_quantized_into( cache["v"][blk_idx], quantize_to_fp4(v_block, self.kv_quant_config), ) else: # Roll cached tokens. cache["k"][:, sink_tokens:sink_tokens + num_rolled_tokens] = \ cache["k"][:, sink_tokens + num_evicted_tokens:sink_tokens + num_evicted_tokens + num_rolled_tokens].clone() cache["v"][:, sink_tokens:sink_tokens + num_rolled_tokens] = \ cache["v"][:, sink_tokens + num_evicted_tokens:sink_tokens + num_evicted_tokens + num_rolled_tokens].clone() # Insert the new key/value tensors. cache["k"][:, local_start_index:local_end_index] = new_k cache["v"][:, local_start_index:local_end_index] = new_v # If a pinned multi-shot sink lives outside position 0, # the rolling shifted everything left by num_evicted_tokens; # pinned_start must follow so it tracks the same data. pinned_shift = update_info.get("pinned_shift", 0) if pinned_shift > 0 and "pinned_start" in cache: cache["pinned_start"].sub_(pinned_shift) elif update_info["action"] == "direct_insert": # Insert directly. local_start_index = update_info["local_start_index"] local_end_index = update_info["local_end_index"] new_k = update_info["new_k"] new_v = update_info["new_v"] if is_quantized: from utils.quant import copy_quantized_into, quantize_to_fp4 blk_sz = int(cache["block_token_size"]) start_blk = local_start_index // blk_sz n_insert_blks = (local_end_index - local_start_index) // blk_sz head_dim = new_k.shape[-1] # iter-26: in-place copy (see note above). for bi in range(n_insert_blks): blk_idx = start_blk + bi ts = bi * blk_sz te = ts + blk_sz k_block = new_k[0, ts:te, :, :].reshape(-1, head_dim).contiguous() v_block = new_v[0, ts:te, :, :].reshape(-1, head_dim).contiguous() copy_quantized_into( cache["k"][blk_idx], quantize_to_fp4(k_block, self.kv_quant_config), ) copy_quantized_into( cache["v"][blk_idx], quantize_to_fp4(v_block, self.kv_quant_config), ) else: # Insert the new key/value tensors. cache["k"][:, local_start_index:local_end_index] = new_k cache["v"][:, local_start_index:local_end_index] = new_v # Update cache indices. kv_cache[block_index]["global_end_index"].fill_(current_end) kv_cache[block_index]["local_end_index"].fill_(local_end_index) def _forward_inference( self, x, t, context, seq_len, clip_fea=None, y=None, kv_cache: dict = None, crossattn_cache: dict = None, current_start: int = 0, cache_start: int = 0, defer_cache_updates: bool = False, ): r""" Run the diffusion model with kv caching. See Algorithm 2 of CausVid paper https://arxiv.org/abs/2412.07772 for details. This function will be run for num_frame times. Process the latent frames one by one (880 tokens each) Args: x (List[Tensor]): List of input video tensors, each with shape [C_in, F, H, W] t (Tensor): Diffusion timesteps tensor of shape [B, F] context (List[Tensor]): List of text embeddings each with shape [L, C] seq_len (`int`): Maximum sequence length for positional encoding clip_fea (Tensor, *optional*): CLIP image features for image-to-video mode y (List[Tensor], *optional*): Conditional video inputs for image-to-video mode, same shape as x Returns: List[Tensor]: List of denoised video tensors with original input shapes [C_out, F, H / 8, W / 8] """ if self.model_type == 'i2v': assert clip_fea is not None and y is not None # params device = self.patch_embedding.weight.device if self.freqs.device != device: self.freqs = self.freqs.to(device) if y is not None: x = [torch.cat([u, v], dim=0) for u, v in zip(x, y)] # embeddings x = [self.patch_embedding(u.unsqueeze(0)) for u in x] # iter-31: publish chunk metadata as Python ints to module-level dict # so deep attention forwards can avoid `.item()` graph breaks. first_shape = tuple(x[0].shape[2:]) _CURRENT_GRID_META["frame_seqlen"] = int(first_shape[1] * first_shape[2]) _CURRENT_GRID_META["num_new_frames"] = int(first_shape[0]) _CURRENT_GRID_META["h"] = int(first_shape[1]) _CURRENT_GRID_META["w"] = int(first_shape[2]) # iter-39 v2: kv_cache scalars (global_end_index, local_end_index, # pinned_start, pinned_len) are published into _CURRENT_GRID_META by # the eager `_call_model` wrapper (utils/wan_5b_wrapper.py) BEFORE # this compiled forward runs. Reading them here via `.item()` would # trigger graph breaks; the wrapper does it in eager Python instead. grid_sizes = torch.stack( [torch.tensor(u.shape[2:], dtype=torch.long) for u in x]) x = [u.flatten(2).transpose(1, 2) for u in x] seq_lens = torch.tensor([u.size(1) for u in x], dtype=torch.long) x = torch.cat(x) # time embeddings if t.dim() == 1: raise NotImplementedError(f"t.shape should be [B, F], but got {t.shape}") bt = t.size(0) t_len = t.size(1) t = t.flatten() e = self.time_embedding( sinusoidal_embedding_1d(self.freq_dim, t).unflatten(0, (bt, t_len)).type_as(x)) e0 = self.time_projection(e).unflatten(2, (6, self.dim)) # B, F, 6, C # context context_lens = None context = self.text_embedding( torch.stack([ torch.cat( [u, u.new_zeros(self.text_len - u.size(0), u.size(1))]) for u in context ])) # arguments kwargs = dict( e=e0, seq_lens=seq_lens, grid_sizes=grid_sizes, freqs=self.freqs, context=context, context_lens=context_lens, block_mask=self.block_mask, t_scale=self.t_scale, use_relative_rope=self.use_relative_rope, method=self.rope_method, original_seq_len=self.original_seq_len, temporal_offset=self.rope_temporal_offset, ) def create_custom_forward(module): def custom_forward(*inputs, **kwargs): return module(*inputs, **kwargs) return custom_forward cache_update_info = None cache_update_infos = [] # Collect cache updates from every block. for block_index, block in enumerate(self.blocks): if torch.is_grad_enabled() and self.gradient_checkpointing: kwargs.update( { "kv_cache": kv_cache[block_index], "current_start": current_start, "cache_start": cache_start } ) result = torch.utils.checkpoint.checkpoint( create_custom_forward(block), x, **kwargs, use_reentrant=False, ) # Handle the result if kv_cache is not None and isinstance(result, tuple): x, block_cache_update_info = result cache_update_infos.append((block_index, block_cache_update_info)) # Keep only basic metadata for later blocks, without the # concrete cache-update payload. cache_update_info = block_cache_update_info[:2] # (current_end, local_end_index) else: x = result else: kwargs.update( { "kv_cache": kv_cache[block_index], "crossattn_cache": crossattn_cache[block_index], "current_start": current_start, "cache_start": cache_start } ) result = block(x, **kwargs) # Handle the result if kv_cache is not None and isinstance(result, tuple): x, block_cache_update_info = result cache_update_infos.append((block_index, block_cache_update_info)) # Keep only basic metadata for later blocks, without the # concrete cache-update payload. cache_update_info = block_cache_update_info[:2] # (current_end, local_end_index) else: x = result # Apply all cache updates after every block has run. For cudagraphs # experiments this can be deferred to the eager wrapper so cache # mutation does not happen inside the compiled forward. if kv_cache is not None and cache_update_infos and not defer_cache_updates: self._apply_cache_updates(kv_cache, cache_update_infos) # head x = self.head(x, e.unsqueeze(2)) # unpatchify x = self.unpatchify(x, grid_sizes) output = torch.stack(x) if kv_cache is not None and defer_cache_updates: return output, cache_update_infos return output def _forward_train( self, x, t, context, seq_len, clean_x=None, aug_t=None, clip_fea=None, y=None, ): r""" Forward pass through the diffusion model Args: x (List[Tensor]): List of input video tensors, each with shape [C_in, F, H, W] t (Tensor): Diffusion timesteps tensor of shape [B] context (List[Tensor]): List of text embeddings each with shape [L, C] seq_len (`int`): Maximum sequence length for positional encoding clip_fea (Tensor, *optional*): CLIP image features for image-to-video mode y (List[Tensor], *optional*): Conditional video inputs for image-to-video mode, same shape as x Returns: List[Tensor]: List of denoised video tensors with original input shapes [C_out, F, H / 8, W / 8] """ if self.model_type == 'i2v': assert clip_fea is not None and y is not None # params device = self.patch_embedding.weight.device if self.freqs.device != device: self.freqs = self.freqs.to(device) # Construct blockwise causal attn mask # Recreate mask when batch size changes to avoid Triton broadcasting bug current_batch_size = x.shape[0] if self.block_mask is None or self._block_mask_batch_size != current_batch_size: self._block_mask_batch_size = current_batch_size if clean_x is not None: if self.independent_first_frame: raise NotImplementedError() else: self.block_mask = self._prepare_teacher_forcing_mask( device, num_frames=x.shape[2], frame_seqlen=x.shape[-2] * x.shape[-1] // (self.patch_size[1] * self.patch_size[2]), num_frame_per_block=self.num_frame_per_block, batch_size=current_batch_size, ) else: if self.independent_first_frame: self.block_mask = self._prepare_blockwise_causal_attn_mask_i2v( device, num_frames=x.shape[2], frame_seqlen=x.shape[-2] * x.shape[-1] // (self.patch_size[1] * self.patch_size[2]), num_frame_per_block=self.num_frame_per_block, batch_size=current_batch_size, ) else: self.block_mask = self._prepare_blockwise_causal_attn_mask( device, num_frames=x.shape[2], frame_seqlen=x.shape[-2] * x.shape[-1] // (self.patch_size[1] * self.patch_size[2]), num_frame_per_block=self.num_frame_per_block, batch_size=current_batch_size, ) if y is not None: x = [torch.cat([u, v], dim=0) for u, v in zip(x, y)] # embeddings x = [self.patch_embedding(u.unsqueeze(0)) for u in x] grid_sizes = torch.stack( [torch.tensor(u.shape[2:], dtype=torch.long) for u in x]) x = [u.flatten(2).transpose(1, 2) for u in x] seq_lens = torch.tensor([u.size(1) for u in x], dtype=torch.long) max_len = int(seq_lens.max().item()) assert max_len > 0, "Token sequence length is zero after patch embedding" # Pad all samples to the batch max length instead of the first sample length x = torch.cat([ torch.cat([u, u.new_zeros(1, max_len - u.size(1), u.size(2))], dim=1) for u in x ]) # time embeddings if t.dim() == 1: raise NotImplementedError(f"t.shape should be [B, F], but got {t.shape}") bt = t.size(0) t_len = t.size(1) t_ori_shape = t.shape t = t.flatten() e = self.time_embedding( sinusoidal_embedding_1d(self.freq_dim, t).unflatten(0, (bt, t_len)).type_as(x)) e0 = self.time_projection(e).unflatten(2, (6, self.dim)) # B, F, 6, C # context context_lens = None context = self.text_embedding( torch.stack([ torch.cat( [u, u.new_zeros(self.text_len - u.size(0), u.size(1))]) for u in context ])) if clean_x is not None: clean_x = [self.patch_embedding(u.unsqueeze(0)) for u in clean_x] clean_x = [u.flatten(2).transpose(1, 2) for u in clean_x] seq_lens_clean = torch.tensor([u.size(1) for u in clean_x], dtype=torch.long) clean_x = torch.cat([ torch.cat([u, u.new_zeros(1, seq_lens_clean[0] - u.size(1), u.size(2))], dim=1) for u in clean_x ]) x = torch.cat([clean_x, x], dim=1) if aug_t is None: aug_t = torch.zeros(t_ori_shape, device=t.device, dtype=t.dtype) bt_clean = aug_t.size(0) t_clean_len = aug_t.size(1) aug_t = aug_t.flatten() e_clean = self.time_embedding( sinusoidal_embedding_1d(self.freq_dim, aug_t).unflatten(0, (bt_clean, t_clean_len)).type_as(x)) e0_clean = self.time_projection(e_clean).unflatten(2, (6, self.dim)) e0 = torch.cat([e0_clean, e0], dim=1) # arguments kwargs = dict( e=e0, seq_lens=seq_lens, grid_sizes=grid_sizes, freqs=self.freqs, context=context, context_lens=context_lens, block_mask=self.block_mask, t_scale=self.t_scale, method=self.rope_method, original_seq_len=self.original_seq_len, temporal_offset=self.rope_temporal_offset, ) def create_custom_forward(module): def custom_forward(*inputs, **kwargs): return module(*inputs, **kwargs) return custom_forward for block in self.blocks: if torch.is_grad_enabled() and self.gradient_checkpointing: x = torch.utils.checkpoint.checkpoint( create_custom_forward(block), x, **kwargs, use_reentrant=False, ) else: x = block(x, **kwargs) if clean_x is not None: x = x[:, x.shape[1] // 2:] # head x = self.head(x, e.unsqueeze(2)) # unpatchify x = self.unpatchify(x, grid_sizes) return torch.stack(x) def forward( self, *args, **kwargs ): if kwargs.get('kv_cache', None) is not None: return self._forward_inference(*args, **kwargs) else: return self._forward_train(*args, **kwargs) def unpatchify(self, x, grid_sizes): r""" Reconstruct video tensors from patch embeddings. Args: x (List[Tensor]): List of patchified features, each with shape [L, C_out * prod(patch_size)] grid_sizes (Tensor): Original spatial-temporal grid dimensions before patching, shape [B, 3] (3 dimensions correspond to F_patches, H_patches, W_patches) Returns: List[Tensor]: Reconstructed video tensors with shape [C_out, F, H / 8, W / 8] """ c = self.out_dim out = [] for u, v in zip(x, grid_sizes.tolist()): u = u[:math.prod(v)].view(*v, *self.patch_size, c) u = torch.einsum('fhwpqrc->cfphqwr', u) u = u.reshape(c, *[i * j for i, j in zip(v, self.patch_size)]) out.append(u) return out def init_weights(self): r""" Initialize model parameters using Xavier initialization. """ # basic init for m in self.modules(): if isinstance(m, nn.Linear): nn.init.xavier_uniform_(m.weight) if m.bias is not None: nn.init.zeros_(m.bias) # init embeddings nn.init.xavier_uniform_(self.patch_embedding.weight.flatten(1)) for m in self.text_embedding.modules(): if isinstance(m, nn.Linear): nn.init.normal_(m.weight, std=.02) for m in self.time_embedding.modules(): if isinstance(m, nn.Linear): nn.init.normal_(m.weight, std=.02) # init output layer nn.init.zeros_(self.head.head.weight)