Upload folder using huggingface_hub

transformer/.ipynb_checkpoints/config-checkpoint.json

@@ -0,0 +1,27 @@
+{
+  "_class_name": "OmniGen2Transformer2DModel",
+  "_diffusers_version": "0.33.1",
+  "axes_dim_rope": [
+    40,
+    40,
+    40
+  ],
+  "axes_lens": [
+    1024,
+    1664,
+    1664
+  ],
+  "ffn_dim_multiplier": null,
+  "hidden_size": 2520,
+  "in_channels": 16,
+  "multiple_of": 256,
+  "norm_eps": 1e-05,
+  "num_attention_heads": 21,
+  "num_kv_heads": 7,
+  "num_layers": 32,
+  "num_refiner_layers": 2,
+  "out_channels": null,
+  "patch_size": 2,
+  "text_feat_dim": 2048,
+  "timestep_scale": 1000.0
+}

transformer/config.json

@@ -0,0 +1,27 @@
+{
+  "_class_name": "OmniGen2Transformer2DModel",
+  "_diffusers_version": "0.33.1",
+  "axes_dim_rope": [
+    40,
+    40,
+    40
+  ],
+  "axes_lens": [
+    1024,
+    1664,
+    1664
+  ],
+  "ffn_dim_multiplier": null,
+  "hidden_size": 2520,
+  "in_channels": 16,
+  "multiple_of": 256,
+  "norm_eps": 1e-05,
+  "num_attention_heads": 21,
+  "num_kv_heads": 7,
+  "num_layers": 32,
+  "num_refiner_layers": 2,
+  "out_channels": null,
+  "patch_size": 2,
+  "text_feat_dim": 2048,
+  "timestep_scale": 1000.0
+}

transformer/diffusion_pytorch_model.safetensors

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

transformer/transformer_omnigen2.py

@@ -0,0 +1,2104 @@
+import warnings
+import itertools
+from typing import Any, Dict, List, Optional, Tuple, Union
+import math
+
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+from einops import rearrange, repeat
+
+from diffusers.configuration_utils import ConfigMixin, register_to_config
+from diffusers.loaders import PeftAdapterMixin
+from diffusers.loaders.single_file_model import FromOriginalModelMixin
+from diffusers.utils import USE_PEFT_BACKEND, logging, scale_lora_layers, unscale_lora_layers
+from diffusers.models.attention_processor import Attention
+from diffusers.models.modeling_outputs import Transformer2DModelOutput
+from diffusers.models.modeling_utils import ModelMixin
+from diffusers.models.embeddings import get_1d_rotary_pos_embed
+from diffusers.models.activations import get_activation
+from diffusers.models.embeddings import Timesteps
+
+import importlib.util
+import sys
+
+# The package importlib_metadata is in a different place, depending on the python version.
+if sys.version_info < (3, 8):
+    import importlib_metadata
+else:
+    import importlib.metadata as importlib_metadata
+
+def _is_package_available(pkg_name: str):
+    pkg_exists = importlib.util.find_spec(pkg_name) is not None
+    pkg_version = "N/A"
+
+    if pkg_exists:
+        try:
+            pkg_version = importlib_metadata.version(pkg_name)
+        except (ImportError, importlib_metadata.PackageNotFoundError):
+            pkg_exists = False
+
+    return pkg_exists, pkg_version
+
+_triton_available, _triton_version = _is_package_available("triton")
+_flash_attn_available, _flash_attn_version = _is_package_available("flash_attn")
+
+def is_triton_available():
+    return _triton_available
+
+def is_flash_attn_available():
+    return _flash_attn_available
+
+if is_triton_available():
+    # from ...ops.triton.layer_norm import RMSNorm
+    import triton
+    import triton.language as tl
+
+
+    from typing import Callable
+
+
+    def custom_amp_decorator(dec: Callable, cuda_amp_deprecated: bool):
+        def decorator(*args, **kwargs):
+            if cuda_amp_deprecated:
+                kwargs["device_type"] = "cuda"
+            return dec(*args, **kwargs)
+        return decorator
+
+
+    if hasattr(torch.amp, "custom_fwd"): # type: ignore[attr-defined]
+        deprecated = True
+        from torch.amp import custom_fwd, custom_bwd # type: ignore[attr-defined]
+    else:
+        deprecated = False
+        from torch.cuda.amp import custom_fwd, custom_bwd
+
+    custom_fwd = custom_amp_decorator(custom_fwd, deprecated)
+    custom_bwd = custom_amp_decorator(custom_bwd, deprecated)
+
+
+    def triton_autotune_configs():
+        # Return configs with a valid warp count for the current device
+        configs=[]
+        # Maximum threads per block is architecture-dependent in theory, but in reality all are 1024
+        max_threads_per_block=1024
+        # Default to warp size 32 if not defined by device
+        warp_size=getattr(torch.cuda.get_device_properties(torch.cuda.current_device()), "warp_size", 32)
+        # Autotune for warp counts which are powers of 2 and do not exceed thread per block limit
+        warp_count=1
+        while warp_count*warp_size <= max_threads_per_block:
+            configs.append(triton.Config({}, num_warps=warp_count))
+            warp_count*=2
+        return configs
+    
+    @triton.autotune(
+        configs=triton_autotune_configs(),
+        key=["N", "HAS_RESIDUAL", "STORE_RESIDUAL_OUT", "IS_RMS_NORM", "HAS_BIAS"],
+    )
+    # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
+    # @triton.heuristics({"HAS_RESIDUAL": lambda args: args["RESIDUAL"] is not None})
+    @triton.heuristics({"HAS_X1": lambda args: args["X1"] is not None})
+    @triton.heuristics({"HAS_W1": lambda args: args["W1"] is not None})
+    @triton.heuristics({"HAS_B1": lambda args: args["B1"] is not None})
+    @triton.jit
+    def _layer_norm_fwd_1pass_kernel(
+        X,  # pointer to the input
+        Y,  # pointer to the output
+        W,  # pointer to the weights
+        B,  # pointer to the biases
+        RESIDUAL,  # pointer to the residual
+        X1,
+        W1,
+        B1,
+        Y1,
+        RESIDUAL_OUT,  # pointer to the residual
+        ROWSCALE,
+        SEEDS,  # Dropout seeds for each row
+        DROPOUT_MASK,
+        Mean,  # pointer to the mean
+        Rstd,  # pointer to the 1/std
+        stride_x_row,  # how much to increase the pointer when moving by 1 row
+        stride_y_row,
+        stride_res_row,
+        stride_res_out_row,
+        stride_x1_row,
+        stride_y1_row,
+        M,  # number of rows in X
+        N,  # number of columns in X
+        eps,  # epsilon to avoid division by zero
+        dropout_p,  # Dropout probability
+        zero_centered_weight,  # If true, add 1.0 to the weight
+        IS_RMS_NORM: tl.constexpr,
+        BLOCK_N: tl.constexpr,
+        HAS_RESIDUAL: tl.constexpr,
+        STORE_RESIDUAL_OUT: tl.constexpr,
+        HAS_BIAS: tl.constexpr,
+        HAS_DROPOUT: tl.constexpr,
+        STORE_DROPOUT_MASK: tl.constexpr,
+        HAS_ROWSCALE: tl.constexpr,
+        HAS_X1: tl.constexpr,
+        HAS_W1: tl.constexpr,
+        HAS_B1: tl.constexpr,
+    ):
+        # Map the program id to the row of X and Y it should compute.
+        row = tl.program_id(0)
+        X += row * stride_x_row
+        Y += row * stride_y_row
+        if HAS_RESIDUAL:
+            RESIDUAL += row * stride_res_row
+        if STORE_RESIDUAL_OUT:
+            RESIDUAL_OUT += row * stride_res_out_row
+        if HAS_X1:
+            X1 += row * stride_x1_row
+        if HAS_W1:
+            Y1 += row * stride_y1_row
+        # Compute mean and variance
+        cols = tl.arange(0, BLOCK_N)
+        x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
+        if HAS_ROWSCALE:
+            rowscale = tl.load(ROWSCALE + row).to(tl.float32)
+            x *= rowscale
+        if HAS_DROPOUT:
+            # Compute dropout mask
+            # 7 rounds is good enough, and reduces register pressure
+            keep_mask = tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
+            x = tl.where(keep_mask, x / (1.0 - dropout_p), 0.0)
+            if STORE_DROPOUT_MASK:
+                tl.store(DROPOUT_MASK + row * N + cols, keep_mask, mask=cols < N)
+        if HAS_X1:
+            x1 = tl.load(X1 + cols, mask=cols < N, other=0.0).to(tl.float32)
+            if HAS_ROWSCALE:
+                rowscale = tl.load(ROWSCALE + M + row).to(tl.float32)
+                x1 *= rowscale
+            if HAS_DROPOUT:
+                # Compute dropout mask
+                # 7 rounds is good enough, and reduces register pressure
+                keep_mask = (
+                    tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
+                )
+                x1 = tl.where(keep_mask, x1 / (1.0 - dropout_p), 0.0)
+                if STORE_DROPOUT_MASK:
+                    tl.store(DROPOUT_MASK + (M + row) * N + cols, keep_mask, mask=cols < N)
+            x += x1
+        if HAS_RESIDUAL:
+            residual = tl.load(RESIDUAL + cols, mask=cols < N, other=0.0).to(tl.float32)
+            x += residual
+        if STORE_RESIDUAL_OUT:
+            tl.store(RESIDUAL_OUT + cols, x, mask=cols < N)
+        if not IS_RMS_NORM:
+            mean = tl.sum(x, axis=0) / N
+            tl.store(Mean + row, mean)
+            xbar = tl.where(cols < N, x - mean, 0.0)
+            var = tl.sum(xbar * xbar, axis=0) / N
+        else:
+            xbar = tl.where(cols < N, x, 0.0)
+            var = tl.sum(xbar * xbar, axis=0) / N
+        rstd = 1 / tl.sqrt(var + eps)
+        tl.store(Rstd + row, rstd)
+        # Normalize and apply linear transformation
+        mask = cols < N
+        w = tl.load(W + cols, mask=mask).to(tl.float32)
+        if zero_centered_weight:
+            w += 1.0
+        if HAS_BIAS:
+            b = tl.load(B + cols, mask=mask).to(tl.float32)
+        x_hat = (x - mean) * rstd if not IS_RMS_NORM else x * rstd
+        y = x_hat * w + b if HAS_BIAS else x_hat * w
+        # Write output
+        tl.store(Y + cols, y, mask=mask)
+        if HAS_W1:
+            w1 = tl.load(W1 + cols, mask=mask).to(tl.float32)
+            if zero_centered_weight:
+                w1 += 1.0
+            if HAS_B1:
+                b1 = tl.load(B1 + cols, mask=mask).to(tl.float32)
+            y1 = x_hat * w1 + b1 if HAS_B1 else x_hat * w1
+            tl.store(Y1 + cols, y1, mask=mask)
+
+
+    def _layer_norm_fwd(
+        x,
+        weight,
+        bias,
+        eps,
+        residual=None,
+        x1=None,
+        weight1=None,
+        bias1=None,
+        dropout_p=0.0,
+        rowscale=None,
+        out_dtype=None,
+        residual_dtype=None,
+        zero_centered_weight=False,
+        is_rms_norm=False,
+        return_dropout_mask=False,
+        out=None,
+        residual_out=None
+    ):
+        if residual is not None:
+            residual_dtype = residual.dtype
+        M, N = x.shape
+        assert x.stride(-1) == 1
+        if residual is not None:
+            assert residual.stride(-1) == 1
+            assert residual.shape == (M, N)
+        assert weight.shape == (N,)
+        assert weight.stride(-1) == 1
+        if bias is not None:
+            assert bias.stride(-1) == 1
+            assert bias.shape == (N,)
+        if x1 is not None:
+            assert x1.shape == x.shape
+            assert rowscale is None
+            assert x1.stride(-1) == 1
+        if weight1 is not None:
+            assert weight1.shape == (N,)
+            assert weight1.stride(-1) == 1
+        if bias1 is not None:
+            assert bias1.shape == (N,)
+            assert bias1.stride(-1) == 1
+        if rowscale is not None:
+            assert rowscale.is_contiguous()
+            assert rowscale.shape == (M,)
+        # allocate output
+        if out is None:
+            out = torch.empty_like(x, dtype=x.dtype if out_dtype is None else out_dtype)
+        else:
+            assert out.shape == x.shape
+        assert out.stride(-1) == 1
+        if weight1 is not None:
+            y1 = torch.empty_like(out)
+            assert y1.stride(-1) == 1
+        else:
+            y1 = None
+        if (
+            residual is not None
+            or (residual_dtype is not None and residual_dtype != x.dtype)
+            or dropout_p > 0.0
+            or rowscale is not None
+            or x1 is not None
+        ):
+            if residual_out is None:
+                residual_out = torch.empty(
+                    M, N, device=x.device, dtype=residual_dtype if residual_dtype is not None else x.dtype
+                )
+            else:
+                assert residual_out.shape == x.shape
+            assert residual_out.stride(-1) == 1
+        else:
+            residual_out = None
+        mean = torch.empty((M,), dtype=torch.float32, device=x.device) if not is_rms_norm else None
+        rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
+        if dropout_p > 0.0:
+            seeds = torch.randint(
+                2**32, (M if x1 is None else 2 * M,), device=x.device, dtype=torch.int64
+            )
+        else:
+            seeds = None
+        if return_dropout_mask and dropout_p > 0.0:
+            dropout_mask = torch.empty(M if x1 is None else 2 * M, N, device=x.device, dtype=torch.bool)
+        else:
+            dropout_mask = None
+        # Less than 64KB per feature: enqueue fused kernel
+        MAX_FUSED_SIZE = 65536 // x.element_size()
+        BLOCK_N = min(MAX_FUSED_SIZE, triton.next_power_of_2(N))
+        if N > BLOCK_N:
+            raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
+        with torch.cuda.device(x.device.index):
+            _layer_norm_fwd_1pass_kernel[(M,)](
+                x,
+                out,
+                weight,
+                bias,
+                residual,
+                x1,
+                weight1,
+                bias1,
+                y1,
+                residual_out,
+                rowscale,
+                seeds,
+                dropout_mask,
+                mean,
+                rstd,
+                x.stride(0),
+                out.stride(0),
+                residual.stride(0) if residual is not None else 0,
+                residual_out.stride(0) if residual_out is not None else 0,
+                x1.stride(0) if x1 is not None else 0,
+                y1.stride(0) if y1 is not None else 0,
+                M,
+                N,
+                eps,
+                dropout_p,
+                zero_centered_weight,
+                is_rms_norm,
+                BLOCK_N,
+                residual is not None,
+                residual_out is not None,
+                bias is not None,
+                dropout_p > 0.0,
+                dropout_mask is not None,
+                rowscale is not None,
+            )
+        # residual_out is None if residual is None and residual_dtype == input_dtype and dropout_p == 0.0
+        if dropout_mask is not None and x1 is not None:
+            dropout_mask, dropout_mask1 = dropout_mask.tensor_split(2, dim=0)
+        else:
+            dropout_mask1 = None
+        return (
+            out,
+            y1,
+            mean,
+            rstd,
+            residual_out if residual_out is not None else x,
+            seeds,
+            dropout_mask,
+            dropout_mask1,
+        )
+
+    @triton.autotune(
+        configs=triton_autotune_configs(),
+        key=["N", "HAS_DRESIDUAL", "STORE_DRESIDUAL", "IS_RMS_NORM", "HAS_BIAS", "HAS_DROPOUT"],
+    )
+    # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
+    # @triton.heuristics({"HAS_DRESIDUAL": lambda args: args["DRESIDUAL"] is not None})
+    # @triton.heuristics({"STORE_DRESIDUAL": lambda args: args["DRESIDUAL_IN"] is not None})
+    @triton.heuristics({"HAS_ROWSCALE": lambda args: args["ROWSCALE"] is not None})
+    @triton.heuristics({"HAS_DY1": lambda args: args["DY1"] is not None})
+    @triton.heuristics({"HAS_DX1": lambda args: args["DX1"] is not None})
+    @triton.heuristics({"HAS_B1": lambda args: args["DB1"] is not None})
+    @triton.heuristics({"RECOMPUTE_OUTPUT": lambda args: args["Y"] is not None})
+    @triton.jit
+    def _layer_norm_bwd_kernel(
+        X,  # pointer to the input
+        W,  # pointer to the weights
+        B,  # pointer to the biases
+        Y,  # pointer to the output to be recomputed
+        DY,  # pointer to the output gradient
+        DX,  # pointer to the input gradient
+        DW,  # pointer to the partial sum of weights gradient
+        DB,  # pointer to the partial sum of biases gradient
+        DRESIDUAL,
+        W1,
+        DY1,
+        DX1,
+        DW1,
+        DB1,
+        DRESIDUAL_IN,
+        ROWSCALE,
+        SEEDS,
+        Mean,  # pointer to the mean
+        Rstd,  # pointer to the 1/std
+        stride_x_row,  # how much to increase the pointer when moving by 1 row
+        stride_y_row,
+        stride_dy_row,
+        stride_dx_row,
+        stride_dres_row,
+        stride_dy1_row,
+        stride_dx1_row,
+        stride_dres_in_row,
+        M,  # number of rows in X
+        N,  # number of columns in X
+        eps,  # epsilon to avoid division by zero
+        dropout_p,
+        zero_centered_weight,
+        rows_per_program,
+        IS_RMS_NORM: tl.constexpr,
+        BLOCK_N: tl.constexpr,
+        HAS_DRESIDUAL: tl.constexpr,
+        STORE_DRESIDUAL: tl.constexpr,
+        HAS_BIAS: tl.constexpr,
+        HAS_DROPOUT: tl.constexpr,
+        HAS_ROWSCALE: tl.constexpr,
+        HAS_DY1: tl.constexpr,
+        HAS_DX1: tl.constexpr,
+        HAS_B1: tl.constexpr,
+        RECOMPUTE_OUTPUT: tl.constexpr,
+    ):
+        # Map the program id to the elements of X, DX, and DY it should compute.
+        row_block_id = tl.program_id(0)
+        row_start = row_block_id * rows_per_program
+        # Do not early exit if row_start >= M, because we need to write DW and DB
+        cols = tl.arange(0, BLOCK_N)
+        mask = cols < N
+        X += row_start * stride_x_row
+        if HAS_DRESIDUAL:
+            DRESIDUAL += row_start * stride_dres_row
+        if STORE_DRESIDUAL:
+            DRESIDUAL_IN += row_start * stride_dres_in_row
+        DY += row_start * stride_dy_row
+        DX += row_start * stride_dx_row
+        if HAS_DY1:
+            DY1 += row_start * stride_dy1_row
+        if HAS_DX1:
+            DX1 += row_start * stride_dx1_row
+        if RECOMPUTE_OUTPUT:
+            Y += row_start * stride_y_row
+        w = tl.load(W + cols, mask=mask).to(tl.float32)
+        if zero_centered_weight:
+            w += 1.0
+        if RECOMPUTE_OUTPUT and HAS_BIAS:
+            b = tl.load(B + cols, mask=mask, other=0.0).to(tl.float32)
+        if HAS_DY1:
+            w1 = tl.load(W1 + cols, mask=mask).to(tl.float32)
+            if zero_centered_weight:
+                w1 += 1.0
+        dw = tl.zeros((BLOCK_N,), dtype=tl.float32)
+        if HAS_BIAS:
+            db = tl.zeros((BLOCK_N,), dtype=tl.float32)
+        if HAS_DY1:
+            dw1 = tl.zeros((BLOCK_N,), dtype=tl.float32)
+            if HAS_B1:
+                db1 = tl.zeros((BLOCK_N,), dtype=tl.float32)
+        row_end = min((row_block_id + 1) * rows_per_program, M)
+        for row in range(row_start, row_end):
+            # Load data to SRAM
+            x = tl.load(X + cols, mask=mask, other=0).to(tl.float32)
+            dy = tl.load(DY + cols, mask=mask, other=0).to(tl.float32)
+            if HAS_DY1:
+                dy1 = tl.load(DY1 + cols, mask=mask, other=0).to(tl.float32)
+            if not IS_RMS_NORM:
+                mean = tl.load(Mean + row)
+            rstd = tl.load(Rstd + row)
+            # Compute dx
+            xhat = (x - mean) * rstd if not IS_RMS_NORM else x * rstd
+            xhat = tl.where(mask, xhat, 0.0)
+            if RECOMPUTE_OUTPUT:
+                y = xhat * w + b if HAS_BIAS else xhat * w
+                tl.store(Y + cols, y, mask=mask)
+            wdy = w * dy
+            dw += dy * xhat
+            if HAS_BIAS:
+                db += dy
+            if HAS_DY1:
+                wdy += w1 * dy1
+                dw1 += dy1 * xhat
+                if HAS_B1:
+                    db1 += dy1
+            if not IS_RMS_NORM:
+                c1 = tl.sum(xhat * wdy, axis=0) / N
+                c2 = tl.sum(wdy, axis=0) / N
+                dx = (wdy - (xhat * c1 + c2)) * rstd
+            else:
+                c1 = tl.sum(xhat * wdy, axis=0) / N
+                dx = (wdy - xhat * c1) * rstd
+            if HAS_DRESIDUAL:
+                dres = tl.load(DRESIDUAL + cols, mask=mask, other=0).to(tl.float32)
+                dx += dres
+            # Write dx
+            if STORE_DRESIDUAL:
+                tl.store(DRESIDUAL_IN + cols, dx, mask=mask)
+            if HAS_DX1:
+                if HAS_DROPOUT:
+                    keep_mask = (
+                        tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
+                    )
+                    dx1 = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
+                else:
+                    dx1 = dx
+                tl.store(DX1 + cols, dx1, mask=mask)
+            if HAS_DROPOUT:
+                keep_mask = tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
+                dx = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
+            if HAS_ROWSCALE:
+                rowscale = tl.load(ROWSCALE + row).to(tl.float32)
+                dx *= rowscale
+            tl.store(DX + cols, dx, mask=mask)
+
+            X += stride_x_row
+            if HAS_DRESIDUAL:
+                DRESIDUAL += stride_dres_row
+            if STORE_DRESIDUAL:
+                DRESIDUAL_IN += stride_dres_in_row
+            if RECOMPUTE_OUTPUT:
+                Y += stride_y_row
+            DY += stride_dy_row
+            DX += stride_dx_row
+            if HAS_DY1:
+                DY1 += stride_dy1_row
+            if HAS_DX1:
+                DX1 += stride_dx1_row
+        tl.store(DW + row_block_id * N + cols, dw, mask=mask)
+        if HAS_BIAS:
+            tl.store(DB + row_block_id * N + cols, db, mask=mask)
+        if HAS_DY1:
+            tl.store(DW1 + row_block_id * N + cols, dw1, mask=mask)
+            if HAS_B1:
+                tl.store(DB1 + row_block_id * N + cols, db1, mask=mask)
+
+
+    def _layer_norm_bwd(
+        dy,
+        x,
+        weight,
+        bias,
+        eps,
+        mean,
+        rstd,
+        dresidual=None,
+        dy1=None,
+        weight1=None,
+        bias1=None,
+        seeds=None,
+        dropout_p=0.0,
+        rowscale=None,
+        has_residual=False,
+        has_x1=False,
+        zero_centered_weight=False,
+        is_rms_norm=False,
+        x_dtype=None,
+        recompute_output=False,
+    ):
+        M, N = x.shape
+        assert x.stride(-1) == 1
+        assert dy.stride(-1) == 1
+        assert dy.shape == (M, N)
+        if dresidual is not None:
+            assert dresidual.stride(-1) == 1
+            assert dresidual.shape == (M, N)
+        assert weight.shape == (N,)
+        assert weight.stride(-1) == 1
+        if bias is not None:
+            assert bias.stride(-1) == 1
+            assert bias.shape == (N,)
+        if dy1 is not None:
+            assert weight1 is not None
+            assert dy1.shape == dy.shape
+            assert dy1.stride(-1) == 1
+        if weight1 is not None:
+            assert weight1.shape == (N,)
+            assert weight1.stride(-1) == 1
+        if bias1 is not None:
+            assert bias1.shape == (N,)
+            assert bias1.stride(-1) == 1
+        if seeds is not None:
+            assert seeds.is_contiguous()
+            assert seeds.shape == (M if not has_x1 else M * 2,)
+        if rowscale is not None:
+            assert rowscale.is_contiguous()
+            assert rowscale.shape == (M,)
+        # allocate output
+        dx = (
+            torch.empty_like(x)
+            if x_dtype is None
+            else torch.empty(M, N, dtype=x_dtype, device=x.device)
+        )
+        dresidual_in = (
+            torch.empty_like(x)
+            if has_residual
+            and (dx.dtype != x.dtype or dropout_p > 0.0 or rowscale is not None or has_x1)
+            else None
+        )
+        dx1 = torch.empty_like(dx) if (has_x1 and dropout_p > 0.0) else None
+        y = torch.empty(M, N, dtype=dy.dtype, device=dy.device) if recompute_output else None
+        if recompute_output:
+            assert weight1 is None, "recompute_output is not supported with parallel LayerNorm"
+
+        # Less than 64KB per feature: enqueue fused kernel
+        MAX_FUSED_SIZE = 65536 // x.element_size()
+        BLOCK_N = min(MAX_FUSED_SIZE, triton.next_power_of_2(N))
+        if N > BLOCK_N:
+            raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
+        # Increasing the multiple (e.g. 8) will allow more thread blocks to be launched and hide the
+        # latency of the gmem reads/writes, but will increase the time of summing up dw / db.
+        sm_count = torch.cuda.get_device_properties(x.device).multi_processor_count * 8
+        _dw = torch.empty((sm_count, N), dtype=torch.float32, device=weight.device)
+        _db = (
+            torch.empty((sm_count, N), dtype=torch.float32, device=bias.device)
+            if bias is not None
+            else None
+        )
+        _dw1 = torch.empty_like(_dw) if weight1 is not None else None
+        _db1 = torch.empty_like(_db) if bias1 is not None else None
+        rows_per_program = math.ceil(M / sm_count)
+        grid = (sm_count,)
+        with torch.cuda.device(x.device.index):
+            _layer_norm_bwd_kernel[grid](
+                x,
+                weight,
+                bias,
+                y,
+                dy,
+                dx,
+                _dw,
+                _db,
+                dresidual,
+                weight1,
+                dy1,
+                dx1,
+                _dw1,
+                _db1,
+                dresidual_in,
+                rowscale,
+                seeds,
+                mean,
+                rstd,
+                x.stride(0),
+                0 if not recompute_output else y.stride(0),
+                dy.stride(0),
+                dx.stride(0),
+                dresidual.stride(0) if dresidual is not None else 0,
+                dy1.stride(0) if dy1 is not None else 0,
+                dx1.stride(0) if dx1 is not None else 0,
+                dresidual_in.stride(0) if dresidual_in is not None else 0,
+                M,
+                N,
+                eps,
+                dropout_p,
+                zero_centered_weight,
+                rows_per_program,
+                is_rms_norm,
+                BLOCK_N,
+                dresidual is not None,
+                dresidual_in is not None,
+                bias is not None,
+                dropout_p > 0.0,
+            )
+        dw = _dw.sum(0).to(weight.dtype)
+        db = _db.sum(0).to(bias.dtype) if bias is not None else None
+        dw1 = _dw1.sum(0).to(weight1.dtype) if weight1 is not None else None
+        db1 = _db1.sum(0).to(bias1.dtype) if bias1 is not None else None
+        # Don't need to compute dresidual_in separately in this case
+        if has_residual and dx.dtype == x.dtype and dropout_p == 0.0 and rowscale is None:
+            dresidual_in = dx
+        if has_x1 and dropout_p == 0.0:
+            dx1 = dx
+        return (
+            (dx, dw, db, dresidual_in, dx1, dw1, db1)
+            if not recompute_output
+            else (dx, dw, db, dresidual_in, dx1, dw1, db1, y)
+        )
+
+    class LayerNormFn(torch.autograd.Function):
+        @staticmethod
+        def forward(
+            ctx,
+            x,
+            weight,
+            bias,
+            residual=None,
+            x1=None,
+            weight1=None,
+            bias1=None,
+            eps=1e-6,
+            dropout_p=0.0,
+            rowscale=None,
+            prenorm=False,
+            residual_in_fp32=False,
+            zero_centered_weight=False,
+            is_rms_norm=False,
+            return_dropout_mask=False,
+            out=None,
+            residual_out=None
+        ):
+            x_shape_og = x.shape
+            # Check for zero sequence length
+            if x.numel() == 0:
+                ctx.zero_seq_length = True
+                # Only save minimal required tensors for backward
+                # ctx.save_for_backward(weight, bias, weight1, bias1)
+                ctx.x_shape_og = x_shape_og
+                ctx.weight_shape = weight.shape
+                ctx.weight_dtype = weight.dtype
+                ctx.weight_device = weight.device
+
+                ctx.has_bias = bias is not None
+                ctx.bias_shape = bias.shape if bias is not None else None
+                ctx.bias_dtype = bias.dtype if bias is not None else None
+                ctx.bias_device = bias.device if bias is not None else None
+
+                ctx.has_weight1 = weight1 is not None
+                ctx.weight1_shape = weight1.shape if weight1 is not None else None
+                ctx.weight1_dtype = weight1.dtype if weight1 is not None else None
+                ctx.weight1_device = weight1.device if weight1 is not None else None
+
+                ctx.has_bias1 = bias1 is not None
+                ctx.bias1_shape = bias1.shape if bias1 is not None else None
+                ctx.bias1_dtype = bias1.dtype if bias1 is not None else None
+                ctx.bias1_device = bias1.device if bias1 is not None else None
+
+                ctx.has_residual = residual is not None
+                ctx.has_x1 = x1 is not None
+                ctx.dropout_p = dropout_p
+
+                # Handle output tensors with correct dtype
+                y = x  # Preserve input tensor properties
+                y1 = torch.empty_like(x) if x1 is not None else None
+                
+                # Only create residual_out if prenorm is True
+                residual_out = torch.empty(x.shape, 
+                                        dtype=torch.float32 if residual_in_fp32 else x.dtype,
+                                        device=x.device) if prenorm else None
+                
+                # Handle dropout masks
+                dropout_mask = None
+                dropout_mask1 = None
+                if return_dropout_mask:
+                    dropout_mask = torch.empty_like(x, dtype=torch.uint8)
+                    if x1 is not None:
+                        dropout_mask1 = torch.empty_like(x, dtype=torch.uint8)
+
+                # Return based on configuration
+                if not return_dropout_mask:
+                    if weight1 is None:
+                        return y if not prenorm else (y, residual_out)
+                    else:
+                        return (y, y1) if not prenorm else (y, y1, residual_out)
+                else:
+                    if weight1 is None:
+                        return ((y, dropout_mask, dropout_mask1) if not prenorm 
+                            else (y, residual_out, dropout_mask, dropout_mask1))
+                    else:
+                        return ((y, y1, dropout_mask, dropout_mask1) if not prenorm 
+                            else (y, y1, residual_out, dropout_mask, dropout_mask1))
+
+            ctx.zero_seq_length = False  
+            # reshape input data into 2D tensor
+            x = x.reshape(-1, x.shape[-1])
+            if x.stride(-1) != 1:
+                x = x.contiguous()
+            if residual is not None:
+                assert residual.shape == x_shape_og
+                residual = residual.reshape(-1, residual.shape[-1])
+                if residual.stride(-1) != 1:
+                    residual = residual.contiguous()
+            if x1 is not None:
+                assert x1.shape == x_shape_og
+                assert rowscale is None, "rowscale is not supported with parallel LayerNorm"
+                x1 = x1.reshape(-1, x1.shape[-1])
+                if x1.stride(-1) != 1:
+                    x1 = x1.contiguous()
+            weight = weight.contiguous()
+            if bias is not None:
+                bias = bias.contiguous()
+            if weight1 is not None:
+                weight1 = weight1.contiguous()
+            if bias1 is not None:
+                bias1 = bias1.contiguous()
+            if rowscale is not None:
+                rowscale = rowscale.reshape(-1).contiguous()
+            residual_dtype = (
+                residual.dtype
+                if residual is not None
+                else (torch.float32 if residual_in_fp32 else None)
+            )
+            if out is not None:
+                out = out.reshape(-1, out.shape[-1])
+            if residual_out is not None:
+                residual_out = residual_out.reshape(-1, residual_out.shape[-1])
+            y, y1, mean, rstd, residual_out, seeds, dropout_mask, dropout_mask1 = _layer_norm_fwd(
+                x,
+                weight,
+                bias,
+                eps,
+                residual,
+                x1,
+                weight1,
+                bias1,
+                dropout_p=dropout_p,
+                rowscale=rowscale,
+                residual_dtype=residual_dtype,
+                zero_centered_weight=zero_centered_weight,
+                is_rms_norm=is_rms_norm,
+                return_dropout_mask=return_dropout_mask,
+                out=out,
+                residual_out=residual_out
+            )
+            ctx.save_for_backward(
+                residual_out, weight, bias, weight1, bias1, rowscale, seeds, mean, rstd
+            )
+            ctx.x_shape_og = x_shape_og
+            ctx.eps = eps
+            ctx.dropout_p = dropout_p
+            ctx.is_rms_norm = is_rms_norm
+            ctx.has_residual = residual is not None
+            ctx.has_x1 = x1 is not None
+            ctx.prenorm = prenorm
+            ctx.x_dtype = x.dtype
+            ctx.zero_centered_weight = zero_centered_weight
+            y = y.reshape(x_shape_og)
+            y1 = y1.reshape(x_shape_og) if y1 is not None else None
+            residual_out = residual_out.reshape(x_shape_og) if residual_out is not None else None
+            dropout_mask = dropout_mask.reshape(x_shape_og) if dropout_mask is not None else None
+            dropout_mask1 = dropout_mask1.reshape(x_shape_og) if dropout_mask1 is not None else None
+            if not return_dropout_mask:
+                if weight1 is None:
+                    return y if not prenorm else (y, residual_out)
+                else:
+                    return (y, y1) if not prenorm else (y, y1, residual_out)
+            else:
+                if weight1 is None:
+                    return (
+                        (y, dropout_mask, dropout_mask1)
+                        if not prenorm
+                        else (y, residual_out, dropout_mask, dropout_mask1)
+                    )
+                else:
+                    return (
+                        (y, y1, dropout_mask, dropout_mask1)
+                        if not prenorm
+                        else (y, y1, residual_out, dropout_mask, dropout_mask1)
+                    )
+
+        @staticmethod
+        def backward(ctx, dy, *args):
+            if ctx.zero_seq_length:
+                return (
+                    torch.zeros(ctx.x_shape_og, dtype=dy.dtype, device=dy.device),
+                    torch.zeros(ctx.weight_shape, dtype=ctx.weight_dtype, device=ctx.weight_device),
+                    torch.zeros(ctx.bias_shape, dtype=ctx.bias_dtype, device=ctx.bias_device) if ctx.has_bias else None,
+                    torch.zeros(ctx.x_shape_og, dtype=dy.dtype, device=dy.device) if ctx.has_residual else None,
+                    torch.zeros(ctx.x_shape_og, dtype=dy.dtype, device=dy.device) if ctx.has_x1 and ctx.dropout_p > 0.0 else None,
+                    torch.zeros(ctx.weight1_shape, dtype=ctx.weight1_dtype, device=ctx.weight1_device) if ctx.has_weight1 else None,
+                    torch.zeros(ctx.bias1_shape, dtype=ctx.bias1_dtype, device=ctx.bias1_device) if ctx.has_bias1 else None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                    None,
+                )
+            
+            x, weight, bias, weight1, bias1, rowscale, seeds, mean, rstd = ctx.saved_tensors
+            dy = dy.reshape(-1, dy.shape[-1])
+            if dy.stride(-1) != 1:
+                dy = dy.contiguous()
+            assert dy.shape == x.shape
+            if weight1 is not None:
+                dy1, args = args[0], args[1:]
+                dy1 = dy1.reshape(-1, dy1.shape[-1])
+                if dy1.stride(-1) != 1:
+                    dy1 = dy1.contiguous()
+                assert dy1.shape == x.shape
+            else:
+                dy1 = None
+            if ctx.prenorm:
+                dresidual = args[0]
+                dresidual = dresidual.reshape(-1, dresidual.shape[-1])
+                if dresidual.stride(-1) != 1:
+                    dresidual = dresidual.contiguous()
+                assert dresidual.shape == x.shape
+            else:
+                dresidual = None
+            
+            dx, dw, db, dresidual_in, dx1, dw1, db1 = _layer_norm_bwd(
+                dy,
+                x,
+                weight,
+                bias,
+                ctx.eps,
+                mean,
+                rstd,
+                dresidual,
+                dy1,
+                weight1,
+                bias1,
+                seeds,
+                ctx.dropout_p,
+                rowscale,
+                ctx.has_residual,
+                ctx.has_x1,
+                ctx.zero_centered_weight,
+                ctx.is_rms_norm,
+                x_dtype=ctx.x_dtype,
+            )
+            return (
+                dx.reshape(ctx.x_shape_og),
+                dw,
+                db,
+                dresidual_in.reshape(ctx.x_shape_og) if ctx.has_residual else None,
+                dx1.reshape(ctx.x_shape_og) if dx1 is not None else None,
+                dw1,
+                db1,
+                None,
+                None,
+                None,
+                None,
+                None,
+                None,
+                None,
+                None,
+                None,
+                None,
+            )
+        
+    def rms_norm_fn(
+        x,
+        weight,
+        bias,
+        residual=None,
+        x1=None,
+        weight1=None,
+        bias1=None,
+        eps=1e-6,
+        dropout_p=0.0,
+        rowscale=None,
+        prenorm=False,
+        residual_in_fp32=False,
+        zero_centered_weight=False,
+        return_dropout_mask=False,
+        out=None,
+        residual_out=None
+    ):
+        return LayerNormFn.apply(
+            x,
+            weight,
+            bias,
+            residual,
+            x1,
+            weight1,
+            bias1,
+            eps,
+            dropout_p,
+            rowscale,
+            prenorm,
+            residual_in_fp32,
+            zero_centered_weight,
+            True,
+            return_dropout_mask,
+            out,
+            residual_out
+        )
+
+    class RMSNorm(torch.nn.Module):
+        def __init__(self, hidden_size, eps=1e-5, dropout_p=0.0, zero_centered_weight=False,
+                    device=None, dtype=None):
+            factory_kwargs = {"device": device, "dtype": dtype}
+            super().__init__()
+            self.eps = eps
+            if dropout_p > 0.0:
+                self.drop = torch.nn.Dropout(dropout_p)
+            else:
+                self.drop = None
+            self.zero_centered_weight = zero_centered_weight
+            self.weight = torch.nn.Parameter(torch.empty(hidden_size, **factory_kwargs))
+            self.register_parameter("bias", None)
+            self.reset_parameters()
+
+        def reset_parameters(self):
+            if not self.zero_centered_weight:
+                torch.nn.init.ones_(self.weight)
+            else:
+                torch.nn.init.zeros_(self.weight)
+
+        def forward(self, x, residual=None, prenorm=False, residual_in_fp32=False):
+            return rms_norm_fn(
+                x,
+                self.weight,
+                self.bias,
+                residual=residual,
+                eps=self.eps,
+                dropout_p=self.drop.p if self.drop is not None and self.training else 0.0,
+                prenorm=prenorm,
+                residual_in_fp32=residual_in_fp32,
+                zero_centered_weight=self.zero_centered_weight,
+            )
+else:
+    from torch.nn import RMSNorm
+    warnings.warn("Cannot import triton, install triton to use fused RMSNorm for better performance")
+    
+def swiglu(x, y):
+    return F.silu(x.float(), inplace=False).to(x.dtype) * y
+
+logger = logging.get_logger(__name__)
+
+
+class TimestepEmbedding(nn.Module):
+    def __init__(
+        self,
+        in_channels: int,
+        time_embed_dim: int,
+        act_fn: str = "silu",
+        out_dim: int = None,
+        post_act_fn: Optional[str] = None,
+        cond_proj_dim=None,
+        sample_proj_bias=True,
+    ):
+        super().__init__()
+
+        self.linear_1 = nn.Linear(in_channels, time_embed_dim, sample_proj_bias)
+
+        if cond_proj_dim is not None:
+            self.cond_proj = nn.Linear(cond_proj_dim, in_channels, bias=False)
+        else:
+            self.cond_proj = None
+
+        self.act = get_activation(act_fn)
+
+        if out_dim is not None:
+            time_embed_dim_out = out_dim
+        else:
+            time_embed_dim_out = time_embed_dim
+        self.linear_2 = nn.Linear(time_embed_dim, time_embed_dim_out, sample_proj_bias)
+
+        if post_act_fn is None:
+            self.post_act = None
+        else:
+            self.post_act = get_activation(post_act_fn)
+
+        self.initialize_weights()
+        
+    def initialize_weights(self):
+        nn.init.normal_(self.linear_1.weight, std=0.02)
+        nn.init.zeros_(self.linear_1.bias)
+        nn.init.normal_(self.linear_2.weight, std=0.02)
+        nn.init.zeros_(self.linear_2.bias)
+        
+    def forward(self, sample, condition=None):
+        if condition is not None:
+            sample = sample + self.cond_proj(condition)
+        sample = self.linear_1(sample)
+
+        if self.act is not None:
+            sample = self.act(sample)
+
+        sample = self.linear_2(sample)
+
+        if self.post_act is not None:
+            sample = self.post_act(sample)
+        return sample
+    
+def apply_rotary_emb(
+    x: torch.Tensor,
+    freqs_cis: Union[torch.Tensor, Tuple[torch.Tensor]],
+    use_real: bool = True,
+    use_real_unbind_dim: int = -1,
+) -> Tuple[torch.Tensor, torch.Tensor]:
+    """
+    Apply rotary embeddings to input tensors using the given frequency tensor. This function applies rotary embeddings
+    to the given query or key 'x' tensors using the provided frequency tensor 'freqs_cis'. The input tensors are
+    reshaped as complex numbers, and the frequency tensor is reshaped for broadcasting compatibility. The resulting
+    tensors contain rotary embeddings and are returned as real tensors.
+
+    Args:
+        x (`torch.Tensor`):
+            Query or key tensor to apply rotary embeddings. [B, H, S, D] xk (torch.Tensor): Key tensor to apply
+        freqs_cis (`Tuple[torch.Tensor]`): Precomputed frequency tensor for complex exponentials. ([S, D], [S, D],)
+
+    Returns:
+        Tuple[torch.Tensor, torch.Tensor]: Tuple of modified query tensor and key tensor with rotary embeddings.
+    """
+    if use_real:
+        cos, sin = freqs_cis  # [S, D]
+        cos = cos[None, None]
+        sin = sin[None, None]
+        cos, sin = cos.to(x.device), sin.to(x.device)
+
+        if use_real_unbind_dim == -1:
+            # Used for flux, cogvideox, hunyuan-dit
+            x_real, x_imag = x.reshape(*x.shape[:-1], -1, 2).unbind(-1)  # [B, S, H, D//2]
+            x_rotated = torch.stack([-x_imag, x_real], dim=-1).flatten(3)
+        elif use_real_unbind_dim == -2:
+            # Used for Stable Audio, OmniGen and CogView4
+            x_real, x_imag = x.reshape(*x.shape[:-1], 2, -1).unbind(-2)  # [B, S, H, D//2]
+            x_rotated = torch.cat([-x_imag, x_real], dim=-1)
+        else:
+            raise ValueError(f"`use_real_unbind_dim={use_real_unbind_dim}` but should be -1 or -2.")
+
+        out = (x.float() * cos + x_rotated.float() * sin).to(x.dtype)
+
+        return out
+    else:
+        # used for lumina
+        # x_rotated = torch.view_as_complex(x.float().reshape(*x.shape[:-1], -1, 2))
+        x_rotated = torch.view_as_complex(x.float().reshape(*x.shape[:-1], x.shape[-1] // 2, 2))
+        freqs_cis = freqs_cis.unsqueeze(2)
+        x_out = torch.view_as_real(x_rotated * freqs_cis).flatten(3)
+
+        return x_out.type_as(x)
+    
+class OmniGen2RotaryPosEmbed(nn.Module):
+    def __init__(self, theta: int,
+                 axes_dim: Tuple[int, int, int],
+                 axes_lens: Tuple[int, int, int] = (300, 512, 512),
+                 patch_size: int = 2):
+        super().__init__()
+        self.theta = theta
+        self.axes_dim = axes_dim
+        self.axes_lens = axes_lens
+        self.patch_size = patch_size
+
+    @staticmethod
+    def get_freqs_cis(axes_dim: Tuple[int, int, int],
+                      axes_lens: Tuple[int, int, int],
+                      theta: int) -> List[torch.Tensor]:
+        freqs_cis = []
+        freqs_dtype = torch.float32 if torch.backends.mps.is_available() else torch.float64
+        for i, (d, e) in enumerate(zip(axes_dim, axes_lens)):
+            emb = get_1d_rotary_pos_embed(d, e, theta=theta, freqs_dtype=freqs_dtype)
+            freqs_cis.append(emb)
+        return freqs_cis
+
+    def _get_freqs_cis(self, freqs_cis, ids: torch.Tensor) -> torch.Tensor:
+        device = ids.device
+        if ids.device.type == "mps":
+            ids = ids.to("cpu")
+
+        result = []
+        for i in range(len(self.axes_dim)):
+            freqs = freqs_cis[i].to(ids.device)
+            index = ids[:, :, i : i + 1].repeat(1, 1, freqs.shape[-1]).to(torch.int64)
+            result.append(torch.gather(freqs.unsqueeze(0).repeat(index.shape[0], 1, 1), dim=1, index=index))
+        return torch.cat(result, dim=-1).to(device)
+
+    def forward(
+        self,
+        freqs_cis,
+        attention_mask,
+        l_effective_ref_img_len,
+        l_effective_img_len,
+        ref_img_sizes,
+        img_sizes,
+        device
+    ):
+        batch_size = len(attention_mask)
+        p = self.patch_size
+
+        encoder_seq_len = attention_mask.shape[1]
+        l_effective_cap_len = attention_mask.sum(dim=1).tolist()
+
+        seq_lengths = [cap_len + sum(ref_img_len) + img_len for cap_len, ref_img_len, img_len in zip(l_effective_cap_len, l_effective_ref_img_len, l_effective_img_len)]
+
+        max_seq_len = max(seq_lengths)
+        max_ref_img_len = max([sum(ref_img_len) for ref_img_len in l_effective_ref_img_len])
+        max_img_len = max(l_effective_img_len)
+
+        # Create position IDs
+        position_ids = torch.zeros(batch_size, max_seq_len, 3, dtype=torch.int32, device=device)
+
+        for i, (cap_seq_len, seq_len) in enumerate(zip(l_effective_cap_len, seq_lengths)):
+            # add text position ids
+            position_ids[i, :cap_seq_len] = repeat(torch.arange(cap_seq_len, dtype=torch.int32, device=device), "l -> l 3")
+
+            pe_shift = cap_seq_len
+            pe_shift_len = cap_seq_len
+
+            if ref_img_sizes[i] is not None:
+                for ref_img_size, ref_img_len in zip(ref_img_sizes[i], l_effective_ref_img_len[i]):
+                    H, W = ref_img_size
+                    ref_H_tokens, ref_W_tokens = H // p, W // p
+                    assert ref_H_tokens * ref_W_tokens == ref_img_len
+                    # add image position ids
+
+                    row_ids = repeat(torch.arange(ref_H_tokens, dtype=torch.int32, device=device), "h -> h w", w=ref_W_tokens).flatten()
+                    col_ids = repeat(torch.arange(ref_W_tokens, dtype=torch.int32, device=device), "w -> h w", h=ref_H_tokens).flatten()
+                    position_ids[i, pe_shift_len:pe_shift_len + ref_img_len, 0] = pe_shift
+                    position_ids[i, pe_shift_len:pe_shift_len + ref_img_len, 1] = row_ids
+                    position_ids[i, pe_shift_len:pe_shift_len + ref_img_len, 2] = col_ids
+
+                    pe_shift += max(ref_H_tokens, ref_W_tokens)
+                    pe_shift_len += ref_img_len
+
+            H, W = img_sizes[i]
+            H_tokens, W_tokens = H // p, W // p
+            assert H_tokens * W_tokens == l_effective_img_len[i]
+
+            row_ids = repeat(torch.arange(H_tokens, dtype=torch.int32, device=device), "h -> h w", w=W_tokens).flatten()
+            col_ids = repeat(torch.arange(W_tokens, dtype=torch.int32, device=device), "w -> h w", h=H_tokens).flatten()
+
+            assert pe_shift_len + l_effective_img_len[i] == seq_len
+            position_ids[i, pe_shift_len: seq_len, 0] = pe_shift
+            position_ids[i, pe_shift_len: seq_len, 1] = row_ids
+            position_ids[i, pe_shift_len: seq_len, 2] = col_ids
+
+        # Get combined rotary embeddings
+        freqs_cis = self._get_freqs_cis(freqs_cis, position_ids)
+        
+        # create separate rotary embeddings for captions and images
+        cap_freqs_cis = torch.zeros(
+            batch_size, encoder_seq_len, freqs_cis.shape[-1], device=device, dtype=freqs_cis.dtype
+        )
+        ref_img_freqs_cis = torch.zeros(
+            batch_size, max_ref_img_len, freqs_cis.shape[-1], device=device, dtype=freqs_cis.dtype
+        )
+        img_freqs_cis = torch.zeros(
+            batch_size, max_img_len, freqs_cis.shape[-1], device=device, dtype=freqs_cis.dtype
+        )
+
+        for i, (cap_seq_len, ref_img_len, img_len, seq_len) in enumerate(zip(l_effective_cap_len, l_effective_ref_img_len, l_effective_img_len, seq_lengths)):
+            cap_freqs_cis[i, :cap_seq_len] = freqs_cis[i, :cap_seq_len]
+            ref_img_freqs_cis[i, :sum(ref_img_len)] = freqs_cis[i, cap_seq_len:cap_seq_len + sum(ref_img_len)]
+            img_freqs_cis[i, :img_len] = freqs_cis[i, cap_seq_len + sum(ref_img_len):cap_seq_len + sum(ref_img_len) + img_len]
+
+        return (
+            cap_freqs_cis,
+            ref_img_freqs_cis,
+            img_freqs_cis,
+            freqs_cis,
+            l_effective_cap_len,
+            seq_lengths,
+        )
+    
+
+class LuminaRMSNormZero(nn.Module):
+    """
+    Norm layer adaptive RMS normalization zero.
+
+    Parameters:
+        embedding_dim (`int`): The size of each embedding vector.
+    """
+
+    def __init__(
+        self,
+        embedding_dim: int,
+        norm_eps: float,
+        norm_elementwise_affine: bool,
+    ):
+        super().__init__()
+        self.silu = nn.SiLU()
+        self.linear = nn.Linear(
+            min(embedding_dim, 1024),
+            4 * embedding_dim,
+            bias=True,
+        )
+        self.norm = RMSNorm(embedding_dim, eps=norm_eps)
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        emb: Optional[torch.Tensor] = None,
+    ) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
+        emb = self.linear(self.silu(emb))
+        scale_msa, gate_msa, scale_mlp, gate_mlp = emb.chunk(4, dim=1)
+        x = self.norm(x) * (1 + scale_msa[:, None])
+        return x, gate_msa, scale_mlp, gate_mlp
+    
+
+class LuminaLayerNormContinuous(nn.Module):
+    def __init__(
+        self,
+        embedding_dim: int,
+        conditioning_embedding_dim: int,
+        # NOTE: It is a bit weird that the norm layer can be configured to have scale and shift parameters
+        # because the output is immediately scaled and shifted by the projected conditioning embeddings.
+        # Note that AdaLayerNorm does not let the norm layer have scale and shift parameters.
+        # However, this is how it was implemented in the original code, and it's rather likely you should
+        # set `elementwise_affine` to False.
+        elementwise_affine=True,
+        eps=1e-5,
+        bias=True,
+        norm_type="layer_norm",
+        out_dim: Optional[int] = None,
+    ):
+        super().__init__()
+
+        # AdaLN
+        self.silu = nn.SiLU()
+        self.linear_1 = nn.Linear(conditioning_embedding_dim, embedding_dim, bias=bias)
+
+        if norm_type == "layer_norm":
+            self.norm = nn.LayerNorm(embedding_dim, eps, elementwise_affine, bias)
+        elif norm_type == "rms_norm":
+            self.norm = RMSNorm(embedding_dim, eps=eps, elementwise_affine=elementwise_affine)
+        else:
+            raise ValueError(f"unknown norm_type {norm_type}")
+
+        self.linear_2 = None
+        if out_dim is not None:
+            self.linear_2 = nn.Linear(embedding_dim, out_dim, bias=bias)
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        conditioning_embedding: torch.Tensor,
+    ) -> torch.Tensor:
+        # convert back to the original dtype in case `conditioning_embedding`` is upcasted to float32 (needed for hunyuanDiT)
+        emb = self.linear_1(self.silu(conditioning_embedding).to(x.dtype))
+        scale = emb
+        x = self.norm(x) * (1 + scale)[:, None, :]
+
+        if self.linear_2 is not None:
+            x = self.linear_2(x)
+
+        return x
+    
+
+class LuminaFeedForward(nn.Module):
+    r"""
+    A feed-forward layer.
+
+    Parameters:
+        hidden_size (`int`):
+            The dimensionality of the hidden layers in the model. This parameter determines the width of the model's
+            hidden representations.
+        intermediate_size (`int`): The intermediate dimension of the feedforward layer.
+        multiple_of (`int`, *optional*): Value to ensure hidden dimension is a multiple
+            of this value.
+        ffn_dim_multiplier (float, *optional*): Custom multiplier for hidden
+            dimension. Defaults to None.
+    """
+
+    def __init__(
+        self,
+        dim: int,
+        inner_dim: int,
+        multiple_of: Optional[int] = 256,
+        ffn_dim_multiplier: Optional[float] = None,
+    ):
+        super().__init__()
+
+        self.swiglu = swiglu
+        
+        # custom hidden_size factor multiplier
+        if ffn_dim_multiplier is not None:
+            inner_dim = int(ffn_dim_multiplier * inner_dim)
+        inner_dim = multiple_of * ((inner_dim + multiple_of - 1) // multiple_of)
+
+        self.linear_1 = nn.Linear(
+            dim,
+            inner_dim,
+            bias=False,
+        )
+        self.linear_2 = nn.Linear(
+            inner_dim,
+            dim,
+            bias=False,
+        )
+        self.linear_3 = nn.Linear(
+            dim,
+            inner_dim,
+            bias=False,
+        )
+
+    def forward(self, x):
+        h1, h2 = self.linear_1(x), self.linear_3(x)
+        return self.linear_2(self.swiglu(h1, h2))
+
+
+class Lumina2CombinedTimestepCaptionEmbedding(nn.Module):
+    def __init__(
+        self,
+        hidden_size: int = 4096,
+        text_feat_dim: int = 2048,
+        frequency_embedding_size: int = 256,
+        norm_eps: float = 1e-5,
+        timestep_scale: float = 1.0,
+    ) -> None:
+        super().__init__()
+
+        self.time_proj = Timesteps(
+            num_channels=frequency_embedding_size, flip_sin_to_cos=True, downscale_freq_shift=0.0, scale=timestep_scale
+        )
+
+        self.timestep_embedder = TimestepEmbedding(
+            in_channels=frequency_embedding_size, time_embed_dim=min(hidden_size, 1024)
+        )
+
+        self.caption_embedder = nn.Sequential(
+            RMSNorm(text_feat_dim, eps=norm_eps),
+            nn.Linear(text_feat_dim, hidden_size, bias=True),
+        )
+        
+        self._initialize_weights()
+
+    def _initialize_weights(self):
+        nn.init.trunc_normal_(self.caption_embedder[1].weight, std=0.02)
+        nn.init.zeros_(self.caption_embedder[1].bias)
+
+    def forward(
+        self, timestep: torch.Tensor, text_hidden_states: torch.Tensor, dtype: torch.dtype
+    ) -> Tuple[torch.Tensor, torch.Tensor]:
+        timestep_proj = self.time_proj(timestep).to(dtype=dtype)
+        time_embed = self.timestep_embedder(timestep_proj)
+        caption_embed = self.caption_embedder(text_hidden_states)
+        return time_embed, caption_embed
+    
+
+class OmniGen2AttnProcessor:
+    """
+    Processor for implementing scaled dot-product attention.
+    
+    This processor is optimized for PyTorch 2.0 and implements:
+    - Flash attention with variable length sequences
+    - Rotary position embeddings (RoPE)
+    - Query-Key normalization
+    - Proportional attention scaling
+    
+    Args:
+        None
+        
+    Raises:
+        ImportError: If PyTorch version is less than 2.0
+    """
+
+    def __init__(self) -> None:
+        """Initialize the attention processor."""
+        if not hasattr(F, "scaled_dot_product_attention"):
+            raise ImportError(
+                "OmniGen2AttnProcessorFlash2Varlen requires PyTorch 2.0. "
+                "Please upgrade PyTorch to version 2.0 or later."
+            )
+
+    def __call__(
+        self,
+        attn: Attention,
+        hidden_states: torch.Tensor,
+        encoder_hidden_states: torch.Tensor,
+        attention_mask: Optional[torch.Tensor] = None,
+        image_rotary_emb: Optional[torch.Tensor] = None,
+        base_sequence_length: Optional[int] = None,
+    ) -> torch.Tensor:
+        """
+        Process attention computation with flash attention.
+
+        Args:
+            attn: Attention module
+            hidden_states: Hidden states tensor of shape (batch_size, seq_len, hidden_dim)
+            encoder_hidden_states: Encoder hidden states tensor
+            attention_mask: Optional attention mask tensor
+            image_rotary_emb: Optional rotary embeddings for image tokens
+            base_sequence_length: Optional base sequence length for proportional attention
+
+        Returns:
+            torch.Tensor: Processed hidden states after attention computation
+        """
+        batch_size, sequence_length, _ = hidden_states.shape
+
+        # Get Query-Key-Value Pair
+        query = attn.to_q(hidden_states)
+        key = attn.to_k(encoder_hidden_states)
+        value = attn.to_v(encoder_hidden_states)
+
+        query_dim = query.shape[-1]
+        inner_dim = key.shape[-1]
+        head_dim = query_dim // attn.heads
+        dtype = query.dtype
+
+        # Get key-value heads
+        kv_heads = inner_dim // head_dim
+
+        # Reshape tensors for attention computation
+        query = query.view(batch_size, -1, attn.heads, head_dim)
+        key = key.view(batch_size, -1, kv_heads, head_dim)
+        value = value.view(batch_size, -1, kv_heads, head_dim)
+
+        # Apply Query-Key normalization
+        if attn.norm_q is not None:
+            query = attn.norm_q(query)
+        if attn.norm_k is not None:
+            key = attn.norm_k(key)
+
+        # Apply Rotary Position Embeddings
+        if image_rotary_emb is not None:
+            query = apply_rotary_emb(query, image_rotary_emb, use_real=False)
+            key = apply_rotary_emb(key, image_rotary_emb, use_real=False)
+
+        query, key = query.to(dtype), key.to(dtype)
+
+        # Calculate attention scale
+        if base_sequence_length is not None:
+            softmax_scale = math.sqrt(math.log(sequence_length, base_sequence_length)) * attn.scale
+        else:
+            softmax_scale = attn.scale
+
+        # scaled_dot_product_attention expects attention_mask shape to be
+        # (batch, heads, source_length, target_length)
+        if attention_mask is not None:
+            attention_mask = attention_mask.bool().view(batch_size, 1, 1, -1)
+
+        query = query.transpose(1, 2)
+        key = key.transpose(1, 2)
+        value = value.transpose(1, 2)
+        
+        # explicitly repeat key and value to match query length, otherwise using enable_gqa=True results in MATH backend of sdpa in our test of pytorch2.6
+        key = key.repeat_interleave(query.size(-3) // key.size(-3), -3)
+        value = value.repeat_interleave(query.size(-3) // value.size(-3), -3)
+
+        hidden_states = F.scaled_dot_product_attention(
+            query, key, value, attn_mask=attention_mask, scale=softmax_scale
+        )
+        hidden_states = hidden_states.transpose(1, 2).reshape(batch_size, -1, attn.heads * head_dim)
+        hidden_states = hidden_states.type_as(query)
+
+        # Apply output projection
+        hidden_states = attn.to_out[0](hidden_states)
+        hidden_states = attn.to_out[1](hidden_states)
+        
+        return hidden_states
+
+class OmniGen2TransformerBlock(nn.Module):
+    """
+    Transformer block for OmniGen2 model.
+    
+    This block implements a transformer layer with:
+    - Multi-head attention with flash attention
+    - Feed-forward network with SwiGLU activation
+    - RMS normalization
+    - Optional modulation for conditional generation
+    
+    Args:
+        dim: Dimension of the input and output tensors
+        num_attention_heads: Number of attention heads
+        num_kv_heads: Number of key-value heads
+        multiple_of: Multiple of which the hidden dimension should be
+        ffn_dim_multiplier: Multiplier for the feed-forward network dimension
+        norm_eps: Epsilon value for normalization layers
+        modulation: Whether to use modulation for conditional generation
+        use_fused_rms_norm: Whether to use fused RMS normalization
+        use_fused_swiglu: Whether to use fused SwiGLU activation
+    """
+
+    def __init__(
+        self,
+        dim: int,
+        num_attention_heads: int,
+        num_kv_heads: int,
+        multiple_of: int,
+        ffn_dim_multiplier: float,
+        norm_eps: float,
+        modulation: bool = True,
+    ) -> None:
+        """Initialize the transformer block."""
+        super().__init__()
+        self.head_dim = dim // num_attention_heads
+        self.modulation = modulation
+
+        # Initialize attention layer
+        self.attn = Attention(
+            query_dim=dim,
+            cross_attention_dim=None,
+            dim_head=dim // num_attention_heads,
+            qk_norm="rms_norm",
+            heads=num_attention_heads,
+            kv_heads=num_kv_heads,
+            eps=1e-5,
+            bias=False,
+            out_bias=False,
+            processor=OmniGen2AttnProcessor(),
+        )
+
+        # Initialize feed-forward network
+        self.feed_forward = LuminaFeedForward(
+            dim=dim,
+            inner_dim=4 * dim,
+            multiple_of=multiple_of,
+            ffn_dim_multiplier=ffn_dim_multiplier,
+        )
+
+        # Initialize normalization layers
+        if modulation:
+            self.norm1 = LuminaRMSNormZero(
+                embedding_dim=dim,
+                norm_eps=norm_eps,
+                norm_elementwise_affine=True,
+            )
+        else:
+            self.norm1 = RMSNorm(dim, eps=norm_eps)
+
+        self.ffn_norm1 = RMSNorm(dim, eps=norm_eps)
+        self.norm2 = RMSNorm(dim, eps=norm_eps)
+        self.ffn_norm2 = RMSNorm(dim, eps=norm_eps)
+
+        self.initialize_weights()
+
+    def initialize_weights(self) -> None:
+        """
+        Initialize the weights of the transformer block.
+        
+        Uses Xavier uniform initialization for linear layers and zero initialization for biases.
+        """
+        nn.init.xavier_uniform_(self.attn.to_q.weight)
+        nn.init.xavier_uniform_(self.attn.to_k.weight)
+        nn.init.xavier_uniform_(self.attn.to_v.weight)
+        nn.init.xavier_uniform_(self.attn.to_out[0].weight)
+
+        nn.init.xavier_uniform_(self.feed_forward.linear_1.weight)
+        nn.init.xavier_uniform_(self.feed_forward.linear_2.weight)
+        nn.init.xavier_uniform_(self.feed_forward.linear_3.weight)
+        
+        if self.modulation:
+            nn.init.zeros_(self.norm1.linear.weight)
+            nn.init.zeros_(self.norm1.linear.bias)
+
+    def forward(
+        self,
+        hidden_states: torch.Tensor,
+        attention_mask: torch.Tensor,
+        image_rotary_emb: torch.Tensor,
+        temb: Optional[torch.Tensor] = None,
+    ) -> torch.Tensor:
+        """
+        Forward pass of the transformer block.
+
+        Args:
+            hidden_states: Input hidden states tensor
+            attention_mask: Attention mask tensor
+            image_rotary_emb: Rotary embeddings for image tokens
+            temb: Optional timestep embedding tensor
+
+        Returns:
+            torch.Tensor: Output hidden states after transformer block processing
+        """
+        if self.modulation:
+            if temb is None:
+                raise ValueError("temb must be provided when modulation is enabled")
+                
+            norm_hidden_states, gate_msa, scale_mlp, gate_mlp = self.norm1(hidden_states, temb)
+            attn_output = self.attn(
+                hidden_states=norm_hidden_states,
+                encoder_hidden_states=norm_hidden_states,
+                attention_mask=attention_mask,
+                image_rotary_emb=image_rotary_emb,
+            )
+            hidden_states = hidden_states + gate_msa.unsqueeze(1).tanh() * self.norm2(attn_output)
+            mlp_output = self.feed_forward(self.ffn_norm1(hidden_states) * (1 + scale_mlp.unsqueeze(1)))
+            hidden_states = hidden_states + gate_mlp.unsqueeze(1).tanh() * self.ffn_norm2(mlp_output)
+        else:
+            norm_hidden_states = self.norm1(hidden_states)
+            attn_output = self.attn(
+                hidden_states=norm_hidden_states,
+                encoder_hidden_states=norm_hidden_states,
+                attention_mask=attention_mask,
+                image_rotary_emb=image_rotary_emb,
+            )
+            hidden_states = hidden_states + self.norm2(attn_output)
+            mlp_output = self.feed_forward(self.ffn_norm1(hidden_states))
+            hidden_states = hidden_states + self.ffn_norm2(mlp_output)
+
+        return hidden_states
+
+
+class OmniGen2Transformer2DModel(ModelMixin, ConfigMixin, PeftAdapterMixin, FromOriginalModelMixin):
+    """
+    OmniGen2 Transformer 2D Model.
+    
+    A transformer-based diffusion model for image generation with:
+    - Patch-based image processing
+    - Rotary position embeddings
+    - Multi-head attention
+    - Conditional generation support
+    
+    Args:
+        patch_size: Size of image patches
+        in_channels: Number of input channels
+        out_channels: Number of output channels (defaults to in_channels)
+        hidden_size: Size of hidden layers
+        num_layers: Number of transformer layers
+        num_refiner_layers: Number of refiner layers
+        num_attention_heads: Number of attention heads
+        num_kv_heads: Number of key-value heads
+        multiple_of: Multiple of which the hidden dimension should be
+        ffn_dim_multiplier: Multiplier for feed-forward network dimension
+        norm_eps: Epsilon value for normalization layers
+        axes_dim_rope: Dimensions for rotary position embeddings
+        axes_lens: Lengths for rotary position embeddings
+        text_feat_dim: Dimension of text features
+        timestep_scale: Scale factor for timestep embeddings
+        use_fused_rms_norm: Whether to use fused RMS normalization
+        use_fused_swiglu: Whether to use fused SwiGLU activation
+    """
+
+    _supports_gradient_checkpointing = True
+    _no_split_modules = ["Omnigen2TransformerBlock"]
+    _skip_layerwise_casting_patterns = ["x_embedder", "norm"]
+
+    @register_to_config
+    def __init__(
+        self,
+        patch_size: int = 2,
+        in_channels: int = 16,
+        out_channels: Optional[int] = None,
+        hidden_size: int = 2304,
+        num_layers: int = 26,
+        num_refiner_layers: int = 2,
+        num_attention_heads: int = 24,
+        num_kv_heads: int = 8,
+        multiple_of: int = 256,
+        ffn_dim_multiplier: Optional[float] = None,
+        norm_eps: float = 1e-5,
+        axes_dim_rope: Tuple[int, int, int] = (32, 32, 32),
+        axes_lens: Tuple[int, int, int] = (300, 512, 512),
+        text_feat_dim: int = 1024,
+        timestep_scale: float = 1.0,
+    ) -> None:
+        """Initialize the OmniGen2 transformer model."""
+        super().__init__()
+
+        # Validate configuration
+        if (hidden_size // num_attention_heads) != sum(axes_dim_rope):
+            raise ValueError(
+                f"hidden_size // num_attention_heads ({hidden_size // num_attention_heads}) "
+                f"must equal sum(axes_dim_rope) ({sum(axes_dim_rope)})"
+            )
+        
+        self.out_channels = out_channels or in_channels
+
+        # Initialize embeddings
+        self.rope_embedder = OmniGen2RotaryPosEmbed(
+            theta=10000,
+            axes_dim=axes_dim_rope,
+            axes_lens=axes_lens,
+            patch_size=patch_size,
+        )
+
+        self.x_embedder = nn.Linear(
+            in_features=patch_size * patch_size * in_channels,
+            out_features=hidden_size,
+        )
+
+        self.ref_image_patch_embedder = nn.Linear(
+            in_features=patch_size * patch_size * in_channels,
+            out_features=hidden_size,
+        )
+
+        self.time_caption_embed = Lumina2CombinedTimestepCaptionEmbedding(
+            hidden_size=hidden_size,
+            text_feat_dim=text_feat_dim,
+            norm_eps=norm_eps,
+            timestep_scale=timestep_scale,
+        )
+
+        # Initialize transformer blocks
+        self.noise_refiner = nn.ModuleList([
+            OmniGen2TransformerBlock(
+                hidden_size,
+                num_attention_heads,
+                num_kv_heads,
+                multiple_of,
+                ffn_dim_multiplier,
+                norm_eps,
+                modulation=True,
+            )
+            for _ in range(num_refiner_layers)
+        ])
+
+        self.ref_image_refiner = nn.ModuleList([
+            OmniGen2TransformerBlock(
+                hidden_size,
+                num_attention_heads,
+                num_kv_heads,
+                multiple_of,
+                ffn_dim_multiplier,
+                norm_eps,
+                modulation=True,
+            )
+            for _ in range(num_refiner_layers)
+        ])
+
+        self.context_refiner = nn.ModuleList(
+            [
+                OmniGen2TransformerBlock(
+                    hidden_size,
+                    num_attention_heads,
+                    num_kv_heads,
+                    multiple_of,
+                    ffn_dim_multiplier,
+                    norm_eps,
+                    modulation=False,
+                )
+                for _ in range(num_refiner_layers)
+            ]
+        )
+
+        # 3. Transformer blocks
+        self.layers = nn.ModuleList(
+            [
+                OmniGen2TransformerBlock(
+                    hidden_size,
+                    num_attention_heads,
+                    num_kv_heads,
+                    multiple_of,
+                    ffn_dim_multiplier,
+                    norm_eps,
+                    modulation=True,
+                )
+                for _ in range(num_layers)
+            ]
+        )
+
+        # 4. Output norm & projection
+        self.norm_out = LuminaLayerNormContinuous(
+            embedding_dim=hidden_size,
+            conditioning_embedding_dim=min(hidden_size, 1024),
+            elementwise_affine=False,
+            eps=1e-6,
+            bias=True,
+            out_dim=patch_size * patch_size * self.out_channels,
+        )
+        
+        # Add learnable embeddings to distinguish different images
+        self.image_index_embedding = nn.Parameter(torch.randn(5, hidden_size)) # support max 5 ref images
+
+        self.gradient_checkpointing = False
+
+        self.initialize_weights()
+
+    def initialize_weights(self) -> None:
+        """
+        Initialize the weights of the model.
+        
+        Uses Xavier uniform initialization for linear layers.
+        """
+        nn.init.xavier_uniform_(self.x_embedder.weight)
+        nn.init.constant_(self.x_embedder.bias, 0.0)
+
+        nn.init.xavier_uniform_(self.ref_image_patch_embedder.weight)
+        nn.init.constant_(self.ref_image_patch_embedder.bias, 0.0)
+
+        nn.init.zeros_(self.norm_out.linear_1.weight)
+        nn.init.zeros_(self.norm_out.linear_1.bias)
+        nn.init.zeros_(self.norm_out.linear_2.weight)
+        nn.init.zeros_(self.norm_out.linear_2.bias)
+        
+        nn.init.normal_(self.image_index_embedding, std=0.02)
+
+    def img_patch_embed_and_refine(
+        self,
+        hidden_states,
+        ref_image_hidden_states,
+        padded_img_mask,
+        padded_ref_img_mask,
+        noise_rotary_emb,
+        ref_img_rotary_emb,
+        l_effective_ref_img_len,
+        l_effective_img_len,
+        temb
+    ):
+        batch_size = len(hidden_states)
+        max_combined_img_len = max([img_len + sum(ref_img_len) for img_len, ref_img_len in zip(l_effective_img_len, l_effective_ref_img_len)])
+    
+        hidden_states = self.x_embedder(hidden_states)
+        ref_image_hidden_states = self.ref_image_patch_embedder(ref_image_hidden_states)
+        
+        for i in range(batch_size):
+            shift = 0
+            for j, ref_img_len in enumerate(l_effective_ref_img_len[i]):
+                ref_image_hidden_states[i, shift:shift + ref_img_len, :] = ref_image_hidden_states[i, shift:shift + ref_img_len, :] + self.image_index_embedding[j]
+                shift += ref_img_len
+
+        for layer in self.noise_refiner:
+            hidden_states = layer(hidden_states, padded_img_mask, noise_rotary_emb, temb)
+
+        flat_l_effective_ref_img_len = list(itertools.chain(*l_effective_ref_img_len))
+        num_ref_images = len(flat_l_effective_ref_img_len)
+        max_ref_img_len = max(flat_l_effective_ref_img_len)
+
+        batch_ref_img_mask = ref_image_hidden_states.new_zeros(num_ref_images, max_ref_img_len, dtype=torch.bool)
+        batch_ref_image_hidden_states = ref_image_hidden_states.new_zeros(num_ref_images, max_ref_img_len, self.config.hidden_size)
+        batch_ref_img_rotary_emb = hidden_states.new_zeros(num_ref_images, max_ref_img_len, ref_img_rotary_emb.shape[-1], dtype=ref_img_rotary_emb.dtype)
+        batch_temb = temb.new_zeros(num_ref_images, *temb.shape[1:], dtype=temb.dtype)
+
+        # sequence of ref imgs to batch
+        idx = 0
+        for i in range(batch_size):
+            shift = 0
+            for ref_img_len in l_effective_ref_img_len[i]:
+                batch_ref_img_mask[idx, :ref_img_len] = True
+                batch_ref_image_hidden_states[idx, :ref_img_len] = ref_image_hidden_states[i, shift:shift + ref_img_len]
+                batch_ref_img_rotary_emb[idx, :ref_img_len] = ref_img_rotary_emb[i, shift:shift + ref_img_len]
+                batch_temb[idx] = temb[i]
+                shift += ref_img_len
+                idx += 1
+
+        # refine ref imgs separately
+        for layer in self.ref_image_refiner:
+            batch_ref_image_hidden_states = layer(batch_ref_image_hidden_states, batch_ref_img_mask, batch_ref_img_rotary_emb, batch_temb)
+
+        # batch of ref imgs to sequence
+        idx = 0
+        for i in range(batch_size):
+            shift = 0
+            for ref_img_len in l_effective_ref_img_len[i]:
+                ref_image_hidden_states[i, shift:shift + ref_img_len] = batch_ref_image_hidden_states[idx, :ref_img_len]
+                shift += ref_img_len
+                idx += 1
+            
+        combined_img_hidden_states = hidden_states.new_zeros(batch_size, max_combined_img_len, self.config.hidden_size)
+        for i, (ref_img_len, img_len) in enumerate(zip(l_effective_ref_img_len, l_effective_img_len)):
+            combined_img_hidden_states[i, :sum(ref_img_len)] = ref_image_hidden_states[i, :sum(ref_img_len)]
+            combined_img_hidden_states[i, sum(ref_img_len):sum(ref_img_len) + img_len] = hidden_states[i, :img_len]
+
+        return combined_img_hidden_states
+
+    def flat_and_pad_to_seq(self, hidden_states, ref_image_hidden_states):
+        batch_size = len(hidden_states)
+        p = self.config.patch_size
+        device = hidden_states[0].device
+
+        img_sizes = [(img.size(1), img.size(2)) for img in hidden_states]
+        l_effective_img_len = [(H // p) * (W // p) for (H, W) in img_sizes]
+
+        if ref_image_hidden_states is not None:
+            ref_img_sizes = [[(img.size(1), img.size(2)) for img in imgs] if imgs is not None else None for imgs in ref_image_hidden_states]
+            l_effective_ref_img_len = [[(ref_img_size[0] // p) * (ref_img_size[1] // p) for ref_img_size in _ref_img_sizes] if _ref_img_sizes is not None else [0] for _ref_img_sizes in ref_img_sizes]
+        else:
+            ref_img_sizes = [None for _ in range(batch_size)]
+            l_effective_ref_img_len = [[0] for _ in range(batch_size)]
+
+        max_ref_img_len = max([sum(ref_img_len) for ref_img_len in l_effective_ref_img_len])
+        max_img_len = max(l_effective_img_len)
+
+        # ref image patch embeddings
+        flat_ref_img_hidden_states = []
+        for i in range(batch_size):
+            if ref_img_sizes[i] is not None:
+                imgs = []
+                for ref_img in ref_image_hidden_states[i]:
+                    C, H, W = ref_img.size()
+                    ref_img = rearrange(ref_img, 'c (h p1) (w p2) -> (h w) (p1 p2 c)', p1=p, p2=p)
+                    imgs.append(ref_img)
+
+                img = torch.cat(imgs, dim=0)
+                flat_ref_img_hidden_states.append(img)
+            else:
+                flat_ref_img_hidden_states.append(None)
+
+        # image patch embeddings
+        flat_hidden_states = []
+        for i in range(batch_size):
+            img = hidden_states[i]
+            C, H, W = img.size()
+            
+            img = rearrange(img, 'c (h p1) (w p2) -> (h w) (p1 p2 c)', p1=p, p2=p)
+            flat_hidden_states.append(img)
+        
+        padded_ref_img_hidden_states = torch.zeros(batch_size, max_ref_img_len, flat_hidden_states[0].shape[-1], device=device, dtype=flat_hidden_states[0].dtype)
+        padded_ref_img_mask = torch.zeros(batch_size, max_ref_img_len, dtype=torch.bool, device=device)
+        for i in range(batch_size):
+            if ref_img_sizes[i] is not None:
+                padded_ref_img_hidden_states[i, :sum(l_effective_ref_img_len[i])] = flat_ref_img_hidden_states[i]
+                padded_ref_img_mask[i, :sum(l_effective_ref_img_len[i])] = True
+
+        padded_hidden_states = torch.zeros(batch_size, max_img_len, flat_hidden_states[0].shape[-1], device=device, dtype=flat_hidden_states[0].dtype)
+        padded_img_mask = torch.zeros(batch_size, max_img_len, dtype=torch.bool, device=device)
+        for i in range(batch_size):
+            padded_hidden_states[i, :l_effective_img_len[i]] = flat_hidden_states[i]
+            padded_img_mask[i, :l_effective_img_len[i]] = True
+
+        return (
+            padded_hidden_states,
+            padded_ref_img_hidden_states,
+            padded_img_mask,
+            padded_ref_img_mask,
+            l_effective_ref_img_len,
+            l_effective_img_len,
+            ref_img_sizes,
+            img_sizes,
+        )
+    
+    def forward(
+        self,
+        hidden_states: Union[torch.Tensor, List[torch.Tensor]],
+        timestep: torch.Tensor,
+        text_hidden_states: torch.Tensor,
+        freqs_cis: torch.Tensor,
+        text_attention_mask: torch.Tensor,
+        ref_image_hidden_states: Optional[List[List[torch.Tensor]]] = None,
+        attention_kwargs: Optional[Dict[str, Any]] = None,
+        return_dict: bool = False,
+    ) -> Union[torch.Tensor, Transformer2DModelOutput]:
+        if attention_kwargs is not None:
+            attention_kwargs = attention_kwargs.copy()
+            lora_scale = attention_kwargs.pop("scale", 1.0)
+        else:
+            lora_scale = 1.0
+
+        if USE_PEFT_BACKEND:
+            # weight the lora layers by setting `lora_scale` for each PEFT layer
+            scale_lora_layers(self, lora_scale)
+        else:
+            if attention_kwargs is not None and attention_kwargs.get("scale", None) is not None:
+                logger.warning(
+                    "Passing `scale` via `attention_kwargs` when not using the PEFT backend is ineffective."
+                )
+
+        # 1. Condition, positional & patch embedding
+        batch_size = len(hidden_states)
+        is_hidden_states_tensor = isinstance(hidden_states, torch.Tensor)
+
+        if is_hidden_states_tensor:
+            assert hidden_states.ndim == 4
+            hidden_states = [_hidden_states for _hidden_states in hidden_states]
+
+        device = hidden_states[0].device
+
+        temb, text_hidden_states = self.time_caption_embed(timestep, text_hidden_states, hidden_states[0].dtype)
+
+        (
+            hidden_states,
+            ref_image_hidden_states,
+            img_mask,
+            ref_img_mask,
+            l_effective_ref_img_len,
+            l_effective_img_len,
+            ref_img_sizes,
+            img_sizes,
+        ) = self.flat_and_pad_to_seq(hidden_states, ref_image_hidden_states)
+        
+        (
+            context_rotary_emb,
+            ref_img_rotary_emb,
+            noise_rotary_emb,
+            rotary_emb,
+            encoder_seq_lengths,
+            seq_lengths,
+        ) = self.rope_embedder(
+            freqs_cis,
+            text_attention_mask,
+            l_effective_ref_img_len,
+            l_effective_img_len,
+            ref_img_sizes,
+            img_sizes,
+            device,
+        )
+
+        # 2. Context refinement
+        for layer in self.context_refiner:
+            text_hidden_states = layer(text_hidden_states, text_attention_mask, context_rotary_emb)
+        
+        combined_img_hidden_states = self.img_patch_embed_and_refine(
+            hidden_states,
+            ref_image_hidden_states,
+            img_mask,
+            ref_img_mask,
+            noise_rotary_emb,
+            ref_img_rotary_emb,
+            l_effective_ref_img_len,
+            l_effective_img_len,
+            temb,
+        )
+
+        # 3. Joint Transformer blocks
+        max_seq_len = max(seq_lengths)
+
+        attention_mask = hidden_states.new_zeros(batch_size, max_seq_len, dtype=torch.bool)
+        joint_hidden_states = hidden_states.new_zeros(batch_size, max_seq_len, self.config.hidden_size)
+        for i, (encoder_seq_len, seq_len) in enumerate(zip(encoder_seq_lengths, seq_lengths)):
+            attention_mask[i, :seq_len] = True
+            joint_hidden_states[i, :encoder_seq_len] = text_hidden_states[i, :encoder_seq_len]
+            joint_hidden_states[i, encoder_seq_len:seq_len] = combined_img_hidden_states[i, :seq_len - encoder_seq_len]
+
+        hidden_states = joint_hidden_states
+
+        for layer_idx, layer in enumerate(self.layers):
+            if torch.is_grad_enabled() and self.gradient_checkpointing:
+                hidden_states = self._gradient_checkpointing_func(
+                    layer, hidden_states, attention_mask, rotary_emb, temb
+                )
+            else:
+                hidden_states = layer(hidden_states, attention_mask, rotary_emb, temb)
+
+        # 4. Output norm & projection
+        hidden_states = self.norm_out(hidden_states, temb)
+
+        p = self.config.patch_size
+        output = []
+        for i, (img_size, img_len, seq_len) in enumerate(zip(img_sizes, l_effective_img_len, seq_lengths)):
+            height, width = img_size
+            output.append(rearrange(hidden_states[i][seq_len - img_len:seq_len], '(h w) (p1 p2 c) -> c (h p1) (w p2)', h=height // p, w=width // p, p1=p, p2=p))
+        if is_hidden_states_tensor:
+            output = torch.stack(output, dim=0)
+
+        if USE_PEFT_BACKEND:
+            # remove `lora_scale` from each PEFT layer
+            unscale_lora_layers(self, lora_scale)
+
+        if not return_dict:
+            return output
+        return Transformer2DModelOutput(sample=output)