Kernels:

kernels-community
/

flash-attn-ops

Trusted publisher

Kernel card Files Files and versions

xet

Community

kernels-bot commited on 13 days ago

Commit

481021d

verified ·

1 Parent(s): f59782d

Uploaded using `kernel-builder`.

Browse files

Files changed (5) hide show

build/torch-cuda/_ops.py +1 -1
build/torch-cuda/_ops_compat.py +0 -18
build/torch-cuda/layer_norm.py +139 -70
build/torch-cuda/metadata.json +14 -4
build/torch-cuda/metadata.json.sigstore +1 -1

build/torch-cuda/_ops.py CHANGED Viewed

@@ -22,7 +22,7 @@ def get_backend() -> str:
 def _find_ops_name() -> str:
     kernel_name = "flash_attn_ops"
-    unique_id = "2e4b976"
     backend = get_backend()
     return f"_{kernel_name}_{backend}_{unique_id}"

 def _find_ops_name() -> str:
     kernel_name = "flash_attn_ops"
+    unique_id = "8a730d9"
     backend = get_backend()
     return f"_{kernel_name}_{backend}_{unique_id}"

build/torch-cuda/_ops_compat.py DELETED Viewed

@@ -1,18 +0,0 @@
-"""Compatibility helpers for op namespacing in source and built layouts.
-In the built (Hub) layout, kernel-builder generates an `_ops` module that
-exposes `add_op_namespace_prefix`, which prefixes op names with a unique,
-build-hashed namespace so custom ops never collide across kernels/versions. When
-running directly from source there is no generated `_ops`, so we fall back to a
-fixed namespace.
-"""
-try:
-    from ._ops import add_op_namespace_prefix as _generated_add_op_namespace_prefix
-except ImportError:
-    def _generated_add_op_namespace_prefix(name: str) -> str:
-        return name if "::" in name else f"flash_attn_ops::{name}"
-def add_op_namespace_prefix(name: str) -> str:
-    return _generated_add_op_namespace_prefix(name)

build/torch-cuda/layer_norm.py CHANGED Viewed

@@ -7,18 +7,17 @@
 # The models we train have hidden dim up to 8k anyway (e.g. Llama 70B), so this is fine.
 import math
-from typing import Optional, List
 import torch
 import torch.nn.functional as F
-from torch import Tensor
 import triton
 import triton.language as tl
-from .utils.torch import custom_fwd, custom_bwd
 from .utils.library import triton_op
-from ._ops_compat import add_op_namespace_prefix
 def maybe_contiguous_lastdim(x):
@@ -40,9 +39,17 @@ def triton_autotune_configs():
     warp_size = 32
     if torch.cuda.is_available():
         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
-    return [triton.Config({}, num_warps=warp_count) for warp_count in [1, 2, 4, 8, 16, 32]
-            if warp_count * warp_size <= max_threads_per_block]
     # return [triton.Config({}, num_warps=8)]
@@ -94,9 +101,9 @@ def layer_norm_ref(
         x = x + x1
     if residual is not None:
         x = (x + residual).to(x.dtype)
-    out = F.layer_norm(x.to(weight.dtype), x.shape[-1:], weight=weight, bias=bias, eps=eps).to(
-        dtype
-    )
     if weight1 is None:
         return out if not prenorm else (out, x)
     else:
@@ -155,19 +162,30 @@ def rms_norm_ref(
     if residual is not None:
         x = (x + residual).to(x.dtype)
     rstd = 1 / torch.sqrt((x.square()).mean(dim=-1, keepdim=True) + eps)
-    out = ((x * rstd * weight) + bias if bias is not None else (x * rstd * weight)).to(dtype)
     if weight1 is None:
         return out if not prenorm else (out, x)
     else:
-        out1 = ((x * rstd * weight1) + bias1 if bias1 is not None else (x * rstd * weight1)).to(
-            dtype
-        )
         return (out, out1) if not prenorm else (out, out1, x)
 @triton.autotune(
     configs=triton_autotune_configs(),
-    key=["N", "HAS_RESIDUAL", "STORE_RESIDUAL_OUT", "IS_RMS_NORM", "HAS_BIAS", "HAS_X1", "HAS_W1", "HAS_B1"],
 )
 # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
 # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
@@ -237,7 +255,9 @@ def _layer_norm_fwd_1pass_kernel(
     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)
@@ -250,7 +270,8 @@ def _layer_norm_fwd_1pass_kernel(
             # 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:
@@ -309,7 +330,7 @@ def _layer_norm_fwd(
     is_rms_norm: bool = False,
     return_dropout_mask: bool = False,
     out: Optional[Tensor] = None,
-    residual_out: Optional[Tensor] = None
 ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
     # Need to wrap to handle the case where residual_out is a alias of x, which makes torch.library
     # and torch.compile unhappy. Also allocate memory for out and residual_out if they are None
@@ -355,8 +376,11 @@ def _layer_norm_fwd(
 # [2025-04-28] torch.library.triton_op ignores the schema argument, but here we need the schema
 # since we're returning a tuple of tensors
-@triton_op(add_op_namespace_prefix("layer_norm_fwd_impl"), mutates_args={"out", "residual_out"},
-           schema="(Tensor x, Tensor weight, Tensor bias, float eps, Tensor(a!) out, Tensor? residual, Tensor? x1, Tensor? weight1, Tensor? bias1, float dropout_p, Tensor? rowscale, bool zero_centered_weight, bool is_rms_norm, bool return_dropout_mask, Tensor(a!)? residual_out) -> (Tensor y1, Tensor mean, Tensor rstd, Tensor seeds, Tensor dropout_mask, Tensor dropout_mask1)")
 def _layer_norm_fwd_impl(
     x: Tensor,
     weight: Tensor,
@@ -372,7 +396,7 @@ def _layer_norm_fwd_impl(
     zero_centered_weight: bool = False,
     is_rms_norm: bool = False,
     return_dropout_mask: bool = False,
-    residual_out: Optional[Tensor] = None
 ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
     M, N = x.shape
     assert x.stride(-1) == 1
@@ -407,7 +431,11 @@ def _layer_norm_fwd_impl(
         assert y1.stride(-1) == 1
     else:
         y1 = 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(
@@ -475,7 +503,14 @@ def _layer_norm_fwd_impl(
 @triton.autotune(
     configs=triton_autotune_configs(),
-    key=["N", "HAS_DRESIDUAL", "STORE_DRESIDUAL", "IS_RMS_NORM", "HAS_BIAS", "HAS_DROPOUT"],
 )
 # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
 # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
@@ -609,14 +644,18 @@ def _layer_norm_bwd_kernel(
         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)
@@ -699,11 +738,12 @@ def _layer_norm_bwd(
     return dx, dw, db, dresidual_in, dx1, dw1, db1, y
-@triton_op(add_op_namespace_prefix("layer_norm_bwd_impl"), mutates_args={},
-           schema="(Tensor dy, Tensor x, Tensor weight, Tensor bias, float eps, Tensor mean, Tensor rstd, Tensor? dresidual, Tensor? dy1, Tensor? weight1, Tensor? bias1, Tensor? seeds, float dropout_p, Tensor? rowscale, bool has_residual, bool has_x1, bool zero_centered_weight, bool is_rms_norm, ScalarType? x_dtype, bool recompute_output) -> (Tensor dx, Tensor dw, Tensor db, Tensor dresidual_in, Tensor dx1, Tensor dw1, Tensor db1, Tensor y)",
-           allow_decomposition=False,  # Don't let torch.compile trace inside
-           )
 def _layer_norm_bwd_impl(
     dy: Tensor,
     x: Tensor,
@@ -770,9 +810,15 @@ def _layer_norm_bwd_impl(
         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()
@@ -849,7 +895,6 @@ def _layer_norm_bwd_impl(
 class LayerNormFn(torch.autograd.Function):
     @staticmethod
     def forward(
         ctx,
@@ -870,14 +915,16 @@ class LayerNormFn(torch.autograd.Function):
         return_dropout_mask=False,
         out_dtype=None,
         out=None,
-        residual_out=None
     ):
         x_shape_og = x.shape
         # reshape input data into 2D tensor
         x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
         if residual is not None:
             assert residual.shape == x_shape_og
-            residual = maybe_contiguous_lastdim(residual.reshape(-1, residual.shape[-1]))
         if x1 is not None:
             assert x1.shape == x_shape_og
             assert rowscale is None, "rowscale is not supported with parallel LayerNorm"
@@ -897,24 +944,26 @@ class LayerNormFn(torch.autograd.Function):
             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,
-            out_dtype=out_dtype,
-            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
@@ -930,9 +979,15 @@ class LayerNormFn(torch.autograd.Function):
         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)
@@ -1030,7 +1085,7 @@ def layer_norm_fn(
     return_dropout_mask=False,
     out_dtype=None,
     out=None,
-    residual_out=None
 ):
     return LayerNormFn.apply(
         x,
@@ -1050,7 +1105,7 @@ def layer_norm_fn(
         return_dropout_mask,
         out_dtype,
         out,
-        residual_out
     )
@@ -1071,7 +1126,7 @@ def rms_norm_fn(
     return_dropout_mask=False,
     out_dtype=None,
     out=None,
-    residual_out=None
 ):
     return LayerNormFn.apply(
         x,
@@ -1091,14 +1146,20 @@ def rms_norm_fn(
         return_dropout_mask,
         out_dtype,
         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
@@ -1132,7 +1193,6 @@ class RMSNorm(torch.nn.Module):
 class LayerNormLinearFn(torch.autograd.Function):
     @staticmethod
     @custom_fwd
     def forward(
@@ -1153,7 +1213,9 @@ class LayerNormLinearFn(torch.autograd.Function):
         x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
         if residual is not None:
             assert residual.shape == x_shape_og
-            residual = maybe_contiguous_lastdim(residual.reshape(-1, residual.shape[-1]))
         norm_weight = norm_weight.contiguous()
         norm_bias = maybe_contiguous(norm_bias)
         residual_dtype = (
@@ -1167,17 +1229,23 @@ class LayerNormLinearFn(torch.autograd.Function):
             norm_bias,
             eps,
             residual,
-            out_dtype=None if not torch.is_autocast_enabled() else torch.get_autocast_dtype("cuda"),
             residual_dtype=residual_dtype,
             is_rms_norm=is_rms_norm,
         )
         y = y.reshape(x_shape_og)
-        dtype = torch.get_autocast_dtype("cuda") if torch.is_autocast_enabled() else y.dtype
         linear_weight = linear_weight.to(dtype)
         linear_bias = linear_bias.to(dtype) if linear_bias is not None else None
         out = F.linear(y.to(linear_weight.dtype), linear_weight, linear_bias)
         # We don't store y, will be recomputed in the backward pass to save memory
-        ctx.save_for_backward(residual_out, norm_weight, norm_bias, linear_weight, mean, rstd)
         ctx.x_shape_og = x_shape_og
         ctx.eps = eps
         ctx.is_rms_norm = is_rms_norm
@@ -1198,8 +1266,9 @@ class LayerNormLinearFn(torch.autograd.Function):
         assert dy.shape == x.shape
         if ctx.prenorm:
             dresidual = args[0]
-            dresidual = maybe_contiguous_lastdim(dresidual.reshape(-1, dresidual.shape[-1]))
-            assert dresidual.shape == x.shape
         else:
             dresidual = None
         dx, dnorm_weight, dnorm_bias, dresidual_in, _, _, _, y = _layer_norm_bwd(

 # The models we train have hidden dim up to 8k anyway (e.g. Llama 70B), so this is fine.
 import math
+from typing import List, Optional
 import torch
 import torch.nn.functional as F
 import triton
 import triton.language as tl
+from torch import Tensor
+from ._ops import add_op_namespace_prefix
 from .utils.library import triton_op
+from .utils.torch import custom_bwd, custom_fwd
 def maybe_contiguous_lastdim(x):
     warp_size = 32
     if torch.cuda.is_available():
         warp_size = getattr(torch.cuda.get_device_properties(torch.cuda.current_device()), "warp_size", 32)
+        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
+    return [
+        triton.Config({}, num_warps=warp_count)
+        for warp_count in [1, 2, 4, 8, 16, 32]
+        if warp_count * warp_size <= max_threads_per_block
+    ]
     # return [triton.Config({}, num_warps=8)]
         x = x + x1
     if residual is not None:
         x = (x + residual).to(x.dtype)
+    out = F.layer_norm(
+        x.to(weight.dtype), x.shape[-1:], weight=weight, bias=bias, eps=eps
+    ).to(dtype)
     if weight1 is None:
         return out if not prenorm else (out, x)
     else:
     if residual is not None:
         x = (x + residual).to(x.dtype)
     rstd = 1 / torch.sqrt((x.square()).mean(dim=-1, keepdim=True) + eps)
+    out = ((x * rstd * weight) + bias if bias is not None else (x * rstd * weight)).to(
+        dtype
+    )
     if weight1 is None:
         return out if not prenorm else (out, x)
     else:
+        out1 = (
+            (x * rstd * weight1) + bias1 if bias1 is not None else (x * rstd * weight1)
+        ).to(dtype)
         return (out, out1) if not prenorm else (out, out1, x)
 @triton.autotune(
     configs=triton_autotune_configs(),
+    key=[
+        "N",
+        "HAS_RESIDUAL",
+        "STORE_RESIDUAL_OUT",
+        "IS_RMS_NORM",
+        "HAS_BIAS",
+        "HAS_X1",
+        "HAS_W1",
+        "HAS_B1",
+    ],
 )
 # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
 # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
     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)
             # 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:
     is_rms_norm: bool = False,
     return_dropout_mask: bool = False,
     out: Optional[Tensor] = None,
+    residual_out: Optional[Tensor] = None,
 ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
     # Need to wrap to handle the case where residual_out is a alias of x, which makes torch.library
     # and torch.compile unhappy. Also allocate memory for out and residual_out if they are None
 # [2025-04-28] torch.library.triton_op ignores the schema argument, but here we need the schema
 # since we're returning a tuple of tensors
+@triton_op(
+    add_op_namespace_prefix("layer_norm_fwd_impl"),
+    mutates_args={"out", "residual_out"},
+    schema="(Tensor x, Tensor weight, Tensor bias, float eps, Tensor(a!) out, Tensor? residual, Tensor? x1, Tensor? weight1, Tensor? bias1, float dropout_p, Tensor? rowscale, bool zero_centered_weight, bool is_rms_norm, bool return_dropout_mask, Tensor(a!)? residual_out) -> (Tensor y1, Tensor mean, Tensor rstd, Tensor seeds, Tensor dropout_mask, Tensor dropout_mask1)",
+)
 def _layer_norm_fwd_impl(
     x: Tensor,
     weight: Tensor,
     zero_centered_weight: bool = False,
     is_rms_norm: bool = False,
     return_dropout_mask: bool = False,
+    residual_out: Optional[Tensor] = None,
 ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
     M, N = x.shape
     assert x.stride(-1) == 1
         assert y1.stride(-1) == 1
     else:
         y1 = 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(
 @triton.autotune(
     configs=triton_autotune_configs(),
+    key=[
+        "N",
+        "HAS_DRESIDUAL",
+        "STORE_DRESIDUAL",
+        "IS_RMS_NORM",
+        "HAS_BIAS",
+        "HAS_DROPOUT",
+    ],
 )
 # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
 # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
         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)
     return dx, dw, db, dresidual_in, dx1, dw1, db1, y
+@triton_op(
+    add_op_namespace_prefix("layer_norm_bwd_impl"),
+    mutates_args={},
+    schema="(Tensor dy, Tensor x, Tensor weight, Tensor bias, float eps, Tensor mean, Tensor rstd, Tensor? dresidual, Tensor? dy1, Tensor? weight1, Tensor? bias1, Tensor? seeds, float dropout_p, Tensor? rowscale, bool has_residual, bool has_x1, bool zero_centered_weight, bool is_rms_norm, ScalarType? x_dtype, bool recompute_output) -> (Tensor dx, Tensor dw, Tensor db, Tensor dresidual_in, Tensor dx1, Tensor dw1, Tensor db1, Tensor y)",
+    allow_decomposition=False,  # Don't let torch.compile trace inside
+)
 def _layer_norm_bwd_impl(
     dy: Tensor,
     x: Tensor,
         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()
 class LayerNormFn(torch.autograd.Function):
     @staticmethod
     def forward(
         ctx,
         return_dropout_mask=False,
         out_dtype=None,
         out=None,
+        residual_out=None,
     ):
         x_shape_og = x.shape
         # reshape input data into 2D tensor
         x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
         if residual is not None:
             assert residual.shape == x_shape_og
+            residual = maybe_contiguous_lastdim(
+                residual.reshape(-1, residual.shape[-1])
+            )
         if x1 is not None:
             assert x1.shape == x_shape_og
             assert rowscale is None, "rowscale is not supported with parallel LayerNorm"
             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,
+                out_dtype=out_dtype,
+                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.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)
     return_dropout_mask=False,
     out_dtype=None,
     out=None,
+    residual_out=None,
 ):
     return LayerNormFn.apply(
         x,
         return_dropout_mask,
         out_dtype,
         out,
+        residual_out,
     )
     return_dropout_mask=False,
     out_dtype=None,
     out=None,
+    residual_out=None,
 ):
     return LayerNormFn.apply(
         x,
         return_dropout_mask,
         out_dtype,
         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
 class LayerNormLinearFn(torch.autograd.Function):
     @staticmethod
     @custom_fwd
     def forward(
         x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
         if residual is not None:
             assert residual.shape == x_shape_og
+            residual = maybe_contiguous_lastdim(
+                residual.reshape(-1, residual.shape[-1])
+            )
         norm_weight = norm_weight.contiguous()
         norm_bias = maybe_contiguous(norm_bias)
         residual_dtype = (
             norm_bias,
             eps,
             residual,
+            out_dtype=None
+            if not torch.is_autocast_enabled()
+            else torch.get_autocast_dtype("cuda"),
             residual_dtype=residual_dtype,
             is_rms_norm=is_rms_norm,
         )
         y = y.reshape(x_shape_og)
+        dtype = (
+            torch.get_autocast_dtype("cuda") if torch.is_autocast_enabled() else y.dtype
+        )
         linear_weight = linear_weight.to(dtype)
         linear_bias = linear_bias.to(dtype) if linear_bias is not None else None
         out = F.linear(y.to(linear_weight.dtype), linear_weight, linear_bias)
         # We don't store y, will be recomputed in the backward pass to save memory
+        ctx.save_for_backward(
+            residual_out, norm_weight, norm_bias, linear_weight, mean, rstd
+        )
         ctx.x_shape_og = x_shape_og
         ctx.eps = eps
         ctx.is_rms_norm = is_rms_norm
         assert dy.shape == x.shape
         if ctx.prenorm:
             dresidual = args[0]
+            dresidual = maybe_contiguous_lastdim(
+                dresidual.reshape(-1, dresidual.shape[-1])
+            )
         else:
             dresidual = None
         dx, dnorm_weight, dnorm_bias, dresidual_in, _, _, _, y = _layer_norm_bwd(

build/torch-cuda/metadata.json CHANGED Viewed

@@ -1,6 +1,6 @@
 {
   "name": "flash-attn-ops",
-  "id": "_flash_attn_ops_cuda_2e4b976",
   "version": 1,
   "license": "BSD-3-Clause",
   "python-depends": [],
@@ -11,17 +11,27 @@
     "algorithm": "sha256",
     "files": {
       "__init__.py": "CN2UymE7gbYnzhbAm92qHBHVC/as55EjJsW+EQt5eaM=",
-      "_ops.py": "p2iM6XDGC6V+jZxNrRc7AwioU+OaCDlgkyHbF++T5Vg=",
-      "_ops_compat.py": "xs6BjkzWExaqXt9xniN5UNYqTp0kpEehRyOyx16+a5o=",
       "cross_entropy.py": "y56nmuOtpSNxM7sDZcUSgXKIsIEayfuvrU8bAVog/nw=",
       "flash_attn_ops/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
       "k_activations.py": "+Z3vIyO4JkqBMipKsPvhzmxljtBdIhJCsl/M+/ESqBo=",
-      "layer_norm.py": "BTwZtyGUbjmWsox0rbMomXsc/UaZJAs+4iexOJ2PfIU=",
       "losses.py": "YhTM9q3Qouq8QoFy1hryAXFC2A6scVtnReoNa+l99YQ=",
       "rotary.py": "0gN+YjhUxwL7B1eE/kKS1Kida5rBrjPKs9Z4j4+U8kA=",
       "utils/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
       "utils/library.py": "nKTJmAHNSqc30g95/NPDtvqzP4GWgYN0BNh015AFzYM=",
       "utils/torch.py": "7mVlWFrCGE7XuCLO9oHa5Z8c4HPLYnE0ojYOJHh8JiQ="
     }
   }
 }

 {
   "name": "flash-attn-ops",
+  "id": "_flash_attn_ops_cuda_8a730d9",
   "version": 1,
   "license": "BSD-3-Clause",
   "python-depends": [],
     "algorithm": "sha256",
     "files": {
       "__init__.py": "CN2UymE7gbYnzhbAm92qHBHVC/as55EjJsW+EQt5eaM=",
+      "_ops.py": "qQNowLEFyFPzxnfTsNt9k+zqYB65oQJozdV6LFvAvNQ=",
       "cross_entropy.py": "y56nmuOtpSNxM7sDZcUSgXKIsIEayfuvrU8bAVog/nw=",
       "flash_attn_ops/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
       "k_activations.py": "+Z3vIyO4JkqBMipKsPvhzmxljtBdIhJCsl/M+/ESqBo=",
+      "layer_norm.py": "SVUTedaQF7zdKwg4bNnr8NPWFSUPB+88kSLWyPkykh0=",
       "losses.py": "YhTM9q3Qouq8QoFy1hryAXFC2A6scVtnReoNa+l99YQ=",
       "rotary.py": "0gN+YjhUxwL7B1eE/kKS1Kida5rBrjPKs9Z4j4+U8kA=",
       "utils/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
       "utils/library.py": "nKTJmAHNSqc30g95/NPDtvqzP4GWgYN0BNh015AFzYM=",
       "utils/torch.py": "7mVlWFrCGE7XuCLO9oHa5Z8c4HPLYnE0ojYOJHh8JiQ="
     }
+  },
+  "provenance": {
+    "kernel-builder": {
+      "version": "0.17.0-dev0",
+      "sha": "81580bb92577f2f7228661ca2a221fc052375709",
+      "dirty": false
+    },
+    "kernel": {
+      "sha": "8a730d96c37560ccf1d3e09f7bbccdc886818f33",
+      "dirty": false
+    }
   }
 }

build/torch-cuda/metadata.json.sigstore CHANGED Viewed

@@ -1 +1 @@

- {"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json","verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUOlUfAEbtfthPfGw7fDjT0prT3ggwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwNzAzMTIwOTIwWhcNMjYwNzAzMTIxOTIwWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEDNCxS/AgnBrb/HmNFA2g9P+lJyQ76Lr5fcwlUfxIib61DJ0nvqK5q+sC6kzT3aqQQynTqdudxepSUBjndLyYdKOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUK8gDLuPcRgJgapENEF0jUNA8HTowHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDJlNGI5NzYyMWU5Mzc5NzQxMzJmNTQzZGIyNTExMjY4MDZmOTRiYmUwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMjg2NTg3NTI1MjgvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABnyfiJ5kAAAQDAEcwRQIhANYsKKhwp+jokEG5vWjGgq5Oh0vvovZUd4v2NTT0Wp8WAiBT2YYP+qUfzbHOSI+kOSK8SAFfxrsBuW+9Ny46bCCVAzAKBggqhkjOPQQDAwNnADBkAjASMs+EqJ047/Sw5kd4QT2KFXmJQW9tISAxHvVzUf32uoPWrUyyXztMbTJ4Um6xFKgCMD0+nKl0puOYII6qZOHF5DtmxT6Rjrm0zvh5dMr9ddmrl8WZu1EwKIv0MWBuz3mh0g=="},"tlogEntries":[{"logIndex":"2060197916","logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="},"kindVersion":{"kind":"hashedrekord","version":"0.0.1"},"integratedTime":"1783080561","inclusionPromise":{"signedEntryTimestamp":"MEUCIFrMGIfLhIxLhETxbQ5OfS52T2txLjLdRB8KjTLlODP7AiEAmiPJPKmgyB1/6lv1AaoE+KjykwCJlIrzI0RkDXEc3pc="},"inclusionProof":{"logIndex":"1938293654","rootHash":"gPz2NSUbLz8lIu3LXY9jfQ07DCa5oaOCDWfdLs4tUR8=","treeSize":"1938293730","hashes":["yJha0KKb0usEfZVMpgK6L/B2iAfLVcOksOwrd78VBSU=","+PU3zRAgU24fymk+h38rChHkLnDtSzRexLdLY562loM=","cwtotohNjqIzjkd69PJo95wBbfcVf7JxVFH2xRjaHk8=","dduNWHkPBB7HIn27N/cE/BeBwVyddRaKkOU7z/H684I=","Bz1UNJYdYVUVyd44dg6GziUUnaWQjiDduEH5H3Xc5UQ=","MqGmWuECTxAuHDgw/dWSvGFP+9rG4p1YGKqwfgrY2xQ=","4605als7RalrKpwjXwzLJTrIVGGZe2sDnGmN4LWkUaM=","2O0Ku6E3Gs7WqZeylo/Q1khiKlgJeUpYTIIil+KxG10=","0YjjV5PnMTeytImbzyW9/jPX6jnOolVeuXHyRfr4yng=","b5cJfno78yVRVpk9SUvKzNVf5K8QDlfznvCcHs9Auus=","7t7nfpRXFFwOWiZ7hqJrJVEhGZwD31UptO0/dy5iWfI=","RIZjwQJv61lwM5JOP8iNT/RxKPr522/etVqxVBWn7KA=","b69YlhA5qlDwN70us4bq9ysnikE1ulA7EwKBTb+atY0=","LM7F5iwbk8jG29X4FwpkCN6BqDCfuNQTlz5CyQWt87s=","+/VZ56MsIPxMiyLAodzKXo5TEWdQp36z89qLhpzloAo=","daxmZaajRpZV+JxHiOYZhJBiSKN5ucqjh2WnGbHhirw=","DOCeoSMovIvLExkhIvisow9AuNXgeWs4ECkyR6EcqYU="],"checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n1938293730\ngPz2NSUbLz8lIu3LXY9jfQ07DCa5oaOCDWfdLs4tUR8=\n\n— rekor.sigstore.dev wNI9ajBFAiAzxny9hSItHnCfVkjrF3tB/HaBYDjqUnDn4B9XqDGTtgIhALJIFKr8NsGzLXVHS1OVRwdhUNXSkkTkosfA1AH7Gv6E\n"}},"canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiI2ZGE3Yzg0NmExOTc0M2Q5NDEzYTRmMmNlOGRlYjBiMjM3Y2EzNWQ1MWZkY2U2ZTZkNmQ3YmZjMzdmMDNiNDZhIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FVUNJRFByNmZHVjNaalpxc3lhMUZRTUxFV1BEZHN6Wm5rTk5tYUFwMWtOUGQwbUFpRUEvbC9id0d2ejB4V0VCTDVUcVBnL1h6eXJCWVFzUGZuYUNnL1FyM0tvWVdJPSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlQyeFZaa0ZGWW5SbWRHaFFaa2QzTjJaRWFsUXdjSEpVTTJkbmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDU2UVhwTlZFbDNUMVJKZDFkb1kwNU5hbGwzVG5wQmVrMVVTWGhQVkVsM1YycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZFVGtONFV5OUJaMjVDY21JdlNHMU9Sa0V5WnpsUUsyeEtlVkUzTmt4eU5XWmpkMndLVldaNFNXbGlOakZFU2pCdWRuRkxOWEVyYzBNMmEzcFVNMkZ4VVZGNWJsUnhaSFZrZUdWd1UxVkNhbTVrVEhsWlpFdFBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZMT0dkRUNreDFVR05TWjBwbllYQkZUa1ZHTUdwVlRrRTRTRlJ2ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU5iVlV3V1dwck0wNXFTWGhhVkd0NlRucHJNMDVFUlhwTmJWa3hDazVFVG10WmFra3hUVlJGZVU1cVozZE9iVmsxVGtkS2FWcFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDF0VlRCWmFtc3pUbXBKZUZwVWEzcE9lbXN6VGtSRmVrMXRXVEZPUkU1cldXcEpNVTFVUlhrS1RtcG5kMDV0V1RWT1IwcHBXbFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDF0VlRCWmFtc3pUbXBKZUZwVWEzcE9lbXN6VGtSRmVrMXRXVEVLVGtST2ExbHFTVEZOVkVWNVRtcG5kMDV0V1RWT1IwcHBXbFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUktiRTVIU1RVS1RucFplVTFYVlRWTmVtTTFUbnBSZUUxNlNtMU9WRkY2V2tkSmVVNVVSWGhOYWxrMFRVUmFiVTlVVW1sWmJWVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOYW1jeVRsUm5NMDVVU1RGTmFtZDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtNTVabWxLTld0QlFVRlJSQXBCUldOM1VsRkphRUZPV1hOTFMyaDNjQ3RxYjJ0RlJ6VjJWMnBIWjNFMVQyZ3dkblp2ZGxwVlpEUjJNazVVVkRCWGNEaFhRV2xDVkRKWldWQXJjVlZtQ25waVNFOVRTU3RyVDFOTE9GTkJSbVo0Y25OQ2RWY3JPVTU1TkRaaVEwTldRWHBCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcEJVMDF6SzBVS2NVb3dORGN2VTNjMWEyUTBVVlF5UzBaWWJVcFJWemwwU1ZOQmVFaDJWbnBWWmpNeWRXOVFWM0pWZVhsWWVuUk5ZbFJLTkZWdE5uaEdTMmREVFVRd0t3cHVTMnd3Y0hWUFdVbEpObkZhVDBoR05VUjBiWGhVTmxKcWNtMHdlblpvTldSTmNqbGtaRzF5YkRoWFduVXhSWGRMU1hZd1RWZENkWG96Yldnd1p6MDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}],"timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyjADAgEAMIICwQYJKoZIhvcNAQcCoIICsjCCAq4CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgpjGhppJpk24JdlqEoOPEMsgq+wu9YtJVgdGkOuzyqCICFQCpu8lcYrDDQ12zZZQFZh5Lpez2pRgPMjAyNjA3MDMxMjA5MjBaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHbMIIB1wIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDcwMzEyMDkyMFowLwYJKoZIhvcNAQkEMSIEIHJsNCxDogcjsHrvO50oiTqDvFnaUMtf+LwpYUyV6aUXMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRnMGUCMQDxiFyZHI5+D6bpr9Z7MAiV2vUJ1kJGe5xD4L1SYMm5zG/ywAE1+S2HaWKrQ9otVzYCMFgroTZmzztiLFGlqkKgAbMCckD1i1AP937uB4trvAEAQE+Xf5kW4xUCuPO6RpE1jg=="}]}},"messageSignature":{"messageDigest":{"algorithm":"SHA2_256","digest":"bafIRqGXQ9lBOk8s6N6wsjfKNdUf3Obm1te/w38DtGo="},"signature":"MEUCIDPr6fGV3ZjZqsya1FQMLEWPDdszZnkNNmaAp1kNPd0mAiEA/l/bwGvz0xWEBL5TqPg/XzyrBYQsPfnaCg/Qr3KoYWI="}}

+ {"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json", "verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtGgAwIBAgIUEE03co9plfqEjv6C6fBMqQMVscUwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwNzA4MDkzNjEwWhcNMjYwNzA4MDk0NjEwWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAErXjaDQ6jAg8TSmfpxGnfDbVzV/4Azr1dZRyC9GQGLohqps2ORsMPHzcJ9ddq20aVebDB+ahbgzFBgVFjmU7GnqOCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUQVQu3XRHLv1kLfF81/mJl6UT9yswHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDhhNzMwZDk2YzM3NTYwY2NmMWQzZTA5ZjdiYmNjZGM4ODY4MThmMzMwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMjg5MzEyMDI0MzIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABn0EVuJ4AAAQDAEcwRQIhAK6jlF9Ixz9e/xTCxVOdECF56DzmbJmN76Z45Hw5ub+oAiBhSJOtm/Ss9XDb0T0t7h8cuppT1ke1JG5s8cN8vepl1DAKBggqhkjOPQQDAwNoADBlAjATX4bvuC4htAHAXmnvOhGu7WJaQvFtAfTTdOsuXIbnO8VwxG0pO4Z98JoIUv3a+wQCMQCL13oHjkbPbe3YZy6BN2dfxHMFp15gwfvKlmZ/ana0Ixji6shwQyi+EodOrjS3zBA="}, "tlogEntries":[{"logIndex":"2114062447", "logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="}, "kindVersion":{"kind":"hashedrekord", "version":"0.0.1"}, "integratedTime":"1783503370", "inclusionPromise":{"signedEntryTimestamp":"MEQCIGULcIB/P5bOtSQFL7zECviyDrT1v6HYnNcOSisMKTAVAiAtQ1U+e6QIUw4GftXfOnTSSFvoDjEeaRhy6HC7N1MsGA=="}, "inclusionProof":{"logIndex":"1992158185", "rootHash":"sZ2Cpq9sCXL4mLvge3quWz8dvrjoDma3YudecTfnSfw=", "treeSize":"1992158187", "hashes":["4H9wkoPKjmBb47xG+bIY9itX9zri+FT0XvvdnjsQ2WU=", "BCVjCP0q/u1HhNdfwREn/VV3bL/8RfAaOQuu3qBgDEU=", "kQ7fcc9d76EvAexPLiVHzpMUUy6NzxEE2G8ZvickbcE=", "ZPO+n58XzQZtmoo7R9cVwyUw/rf3ZXHxPwoxNI1AKu0=", "C/83T0LtgCGOI9u8toJmdS2iIcOmXIvOx/gudBzyWWw=", "lwsNHEM6qJIqfbQe7ZvedAABUsLMZLDOD1xWJHQH1UU=", "Usxh6EvyE9A1zFrYsjfJHvqZycxmIXNgyzfacgnlE48=", "VdcdobyFjwqP6NbAvWaSE/pTWrRccgdl/PbLqF1kZWk=", "DKfazLC+UbUHA0cM6HRVM99d45RuRYGI6beeOTbtJi4=", "6WLV5MAiMmlAor5VYea+SrvSf5ZMe66flxXztNFFIhk=", "+OMAiJ79sBwk/ZbFDPilfkgmCem1MbCFpLOfIHSetfc=", "tY8nHkJkYkzR0a8crBzP4XX2ys6vaQR/b0ImUbbsFfk=", "W8OH9B7+IUIWw6Asm1BMB4ZxX27Q2QG2hpA1yZ+WxAQ=", "q0bDCYYM8Kl1AcQlFtAjsMRkDgQDRjrR/EhSNDT6HeI=", "Uh9Q3t0owp3nQij7uUVzswxCf/9lo7BQqoqZwxfV0u4=", "fZbx4AWo+aeMN9Mbraux+rscCGLJEHB7hXu1e/kRgeo=", "D/x3yP7cG/SpQUwSXzpITAOA0UroT/kIiXSF3wR+jT4=", "Bcvbbu+tDL+zr1H2q26c2ATmVu7v/f7JIFweUu7qDK8=", "WajL3ZhY6ZX9meV3VTH9TkrbM31OF4o6+7foL0Uf6PI=", "cyKMKBvHkQtwb2xR5svh+YppANOVdn7X8t4DuKMxICA=", "+/VZ56MsIPxMiyLAodzKXo5TEWdQp36z89qLhpzloAo=", "daxmZaajRpZV+JxHiOYZhJBiSKN5ucqjh2WnGbHhirw=", "DOCeoSMovIvLExkhIvisow9AuNXgeWs4ECkyR6EcqYU="], "checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n1992158187\nsZ2Cpq9sCXL4mLvge3quWz8dvrjoDma3YudecTfnSfw=\n\n— rekor.sigstore.dev wNI9ajBFAiEA6Piq1feQqdbEyoF3CNxMktfdMmkVxfU3oYnhAnvAZ3sCIDfpqxQTqkIBY8rqLC90uIxZqq/qaY3J48VptJw6Dr3Y\n"}}, "canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJjNzUyZGFkZmRlODJhOTQzMzNkNjVjYzJjMTY2OThjOTg2NzFkYzY1ZGUwNzlkNGM1ZGUwNzYyYTM2OGRkOTlkIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FUUNJR2wzSDRnQ1NmVWRnVC95RlpGU3JDRTBFTnh1N3NTdmZNUkFJcEtqNjZmNUFpQXVSMGxvUWxlcmV0a1ZmSTFVSERpZmprS0hVNVJVOFlvZlVWMHNUWFVIRkE9PSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5SSFowRjNTVUpCWjBsVlJVVXdNMk52T1hCc1puRkZhblkyUXpabVFrMXhVVTFXYzJOVmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDU2UVRSTlJHdDZUbXBGZDFkb1kwNU5hbGwzVG5wQk5FMUVhekJPYWtWM1YycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZ5V0dwaFJGRTJha0ZuT0ZSVGJXWndlRWR1WmtSaVZucFdMelJCZW5JeFpGcFNlVU1LT1VkUlIweHZhSEZ3Y3pKUFVuTk5VRWg2WTBvNVpHUnhNakJoVm1WaVJFSXJZV2hpWjNwR1FtZFdSbXB0VlRkSGJuRlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZSVmxGMUNqTllVa2hNZGpGclRHWkdPREV2YlVwc05sVlVPWGx6ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVBSMFV6VFhwQ2EwOVVXbXBOZW1NeFRtcENhbGt5V1hoYVJFNXNDazFFYkcxT01rcHBXVEpPYTFsNlp6Uk9hbWQ0VDBkWmVrMTZRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDlIUlROTmVrSnJUMVJhYWsxNll6Rk9ha0pxV1RKWmVGcEVUbXhOUkd4dFRqSkthVmt5VG1zS1dYcG5ORTVxWjNoUFIxbDZUWHBCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDlIUlROTmVrSnJUMVJhYWsxNll6Rk9ha0pxV1RKWmVGcEVUbXdLVFVSc2JVNHlTbWxaTWs1cldYcG5ORTVxWjNoUFIxbDZUWHBCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUm9hRTU2VFhjS1drUnJNbGw2VFROT1ZGbDNXVEpPYlUxWFVYcGFWRUUxV21wa2FWbHRUbXBhUjAwMFQwUlpORTFVYUcxTmVrMTNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOYW1jMVRYcEZlVTFFU1RCTmVrbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtNHdSVloxU2pSQlFVRlJSQXBCUldOM1VsRkphRUZMTm1wc1JqbEplSG81WlM5NFZFTjRWazlrUlVOR05UWkVlbTFpU20xT056WmFORFZJZHpWMVlpdHZRV2xDYUZOS1QzUnRMMU56Q2psWVJHSXdWREIwTjJnNFkzVndjRlF4YTJVeFNrYzFjemhqVGpoMlpYQnNNVVJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTlCUkVKc1FXcEJWRmcwWW5ZS2RVTTBhSFJCU0VGWWJXNTJUMmhIZFRkWFNtRlJka1owUVdaVVZHUlBjM1ZZU1dKdVR6aFdkM2hITUhCUE5GbzVPRXB2U1ZWMk0yRXJkMUZEVFZGRFRBb3hNMjlJYW10aVVHSmxNMWxhZVRaQ1RqSmtabmhJVFVad01UVm5kMloyUzJ4dFdpOWhibUV3U1hocWFUWnphSGRSZVdrclJXOWtUM0pxVXpONlFrRTlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}], "timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyDADAgEAMIICvwYJKoZIhvcNAQcCoIICsDCCAqwCAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgmbD6vUdbxXI0QN569eetYp4IOQ9dg0vxIhD7gz1ff/sCFA2iJdplolKF1skRWs4ONQT3t1ofGA8yMDI2MDcwODA5MzYxMFowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdowggHWAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwNzA4MDkzNjEwWjAvBgkqhkiG9w0BCQQxIgQgYVOn7Ane8rL/hWMUgnVG/+xhGiVAmD1T68y13JRNSTIwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGYwZAIwVeNtdTaO1QKTryuXnSoZqKm6Ei1zovRzTLROkPmGpc4bHWXt1eNVy/3dj/vPnjOuAjB453HwcuBc5PtsqE+DX3oxvbSt1byD7k1zDs0dAaqO+HpqM4wUKFZG1DThEBdfof4="}]}}, "messageSignature":{"messageDigest":{"algorithm":"SHA2_256", "digest":"x1La396CqUMz1lzCwWaYyYZx3GXeB51MXeB2KjaN2Z0="}, "signature":"MEQCIGl3H4gCSfUdgT/yFZFSrCE0ENxu7sSvfMRAIpKj66f5AiAuR0loQleretkVfI1UHDifjkKHU5RU8YofUV0sTXUHFA=="}}