Hugging Face's logo

Kernels:
kernels-community
/
flash-attn-ops
Trusted publisher

Kernel card Files
xet
 Community
kernels-bot commited on
Commit
7232da3
·
verified ·
1 Parent(s): c2b03f4

Uploaded using `kernel-builder`.

Browse files
Files changed (5) hide show
  1. build/torch-xpu/_ops.py +1 -1
  2. build/torch-xpu/_ops_compat.py +0 -18
  3. build/torch-xpu/layer_norm.py +139 -70
  4. build/torch-xpu/metadata.json +14 -4
  5. build/torch-xpu/metadata.json.sigstore +1 -1
build/torch-xpu/_ops.py CHANGED
@@ -22,7 +22,7 @@ def get_backend() -> str:
22
 
23
  def _find_ops_name() -> str:
24
  kernel_name = "flash_attn_ops"
25
- unique_id = "2e4b976"
26
  backend = get_backend()
27
  return f"_{kernel_name}_{backend}_{unique_id}"
28
 
 
22
 
23
  def _find_ops_name() -> str:
24
  kernel_name = "flash_attn_ops"
25
+ unique_id = "8a730d9"
26
  backend = get_backend()
27
  return f"_{kernel_name}_{backend}_{unique_id}"
28
 
build/torch-xpu/_ops_compat.py DELETED
@@ -1,18 +0,0 @@
1
- """Compatibility helpers for op namespacing in source and built layouts.
2
-
3
- In the built (Hub) layout, kernel-builder generates an `_ops` module that
4
- exposes `add_op_namespace_prefix`, which prefixes op names with a unique,
5
- build-hashed namespace so custom ops never collide across kernels/versions. When
6
- running directly from source there is no generated `_ops`, so we fall back to a
7
- fixed namespace.
8
- """
9
-
10
- try:
11
- from ._ops import add_op_namespace_prefix as _generated_add_op_namespace_prefix
12
- except ImportError:
13
- def _generated_add_op_namespace_prefix(name: str) -> str:
14
- return name if "::" in name else f"flash_attn_ops::{name}"
15
-
16
-
17
- def add_op_namespace_prefix(name: str) -> str:
18
- return _generated_add_op_namespace_prefix(name)
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
build/torch-xpu/layer_norm.py CHANGED
@@ -7,18 +7,17 @@
7
  # The models we train have hidden dim up to 8k anyway (e.g. Llama 70B), so this is fine.
8
 
9
  import math
10
- from typing import Optional, List
11
 
12
  import torch
13
  import torch.nn.functional as F
14
- from torch import Tensor
15
-
16
  import triton
17
  import triton.language as tl
 
18
 
19
- from .utils.torch import custom_fwd, custom_bwd
20
  from .utils.library import triton_op
21
- from ._ops_compat import add_op_namespace_prefix
22
 
23
 
24
  def maybe_contiguous_lastdim(x):
@@ -40,9 +39,17 @@ def triton_autotune_configs():
40
  warp_size = 32
41
  if torch.cuda.is_available():
42
  warp_size = getattr(torch.cuda.get_device_properties(torch.cuda.current_device()), "warp_size", 32)
 
 
 
 
 
43
  # Autotune for warp counts which are powers of 2 and do not exceed thread per block limit
44
- return [triton.Config({}, num_warps=warp_count) for warp_count in [1, 2, 4, 8, 16, 32]
45
- if warp_count * warp_size <= max_threads_per_block]
 
 
 
46
  # return [triton.Config({}, num_warps=8)]
47
 
48
 
@@ -94,9 +101,9 @@ def layer_norm_ref(
94
  x = x + x1
95
  if residual is not None:
96
  x = (x + residual).to(x.dtype)
97
- out = F.layer_norm(x.to(weight.dtype), x.shape[-1:], weight=weight, bias=bias, eps=eps).to(
98
- dtype
99
- )
100
  if weight1 is None:
101
  return out if not prenorm else (out, x)
102
  else:
@@ -155,19 +162,30 @@ def rms_norm_ref(
155
  if residual is not None:
156
  x = (x + residual).to(x.dtype)
157
  rstd = 1 / torch.sqrt((x.square()).mean(dim=-1, keepdim=True) + eps)
158
- out = ((x * rstd * weight) + bias if bias is not None else (x * rstd * weight)).to(dtype)
 
 
159
  if weight1 is None:
160
  return out if not prenorm else (out, x)
161
  else:
162
- out1 = ((x * rstd * weight1) + bias1 if bias1 is not None else (x * rstd * weight1)).to(
163
- dtype
164
- )
165
  return (out, out1) if not prenorm else (out, out1, x)
166
 
167
 
168
  @triton.autotune(
169
  configs=triton_autotune_configs(),
170
- key=["N", "HAS_RESIDUAL", "STORE_RESIDUAL_OUT", "IS_RMS_NORM", "HAS_BIAS", "HAS_X1", "HAS_W1", "HAS_B1"],
 
 
 
 
 
 
 
 
 
171
  )
172
  # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
173
  # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
@@ -237,7 +255,9 @@ def _layer_norm_fwd_1pass_kernel(
237
  if HAS_DROPOUT:
238
  # Compute dropout mask
239
  # 7 rounds is good enough, and reduces register pressure
240
- keep_mask = tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
 
 
241
  x = tl.where(keep_mask, x / (1.0 - dropout_p), 0.0)
242
  if STORE_DROPOUT_MASK:
243
  tl.store(DROPOUT_MASK + row * N + cols, keep_mask, mask=cols < N)
@@ -250,7 +270,8 @@ def _layer_norm_fwd_1pass_kernel(
250
  # Compute dropout mask
251
  # 7 rounds is good enough, and reduces register pressure
252
  keep_mask = (
253
- tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
 
254
  )
255
  x1 = tl.where(keep_mask, x1 / (1.0 - dropout_p), 0.0)
256
  if STORE_DROPOUT_MASK:
@@ -309,7 +330,7 @@ def _layer_norm_fwd(
309
  is_rms_norm: bool = False,
310
  return_dropout_mask: bool = False,
311
  out: Optional[Tensor] = None,
312
- residual_out: Optional[Tensor] = None
313
  ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
314
  # Need to wrap to handle the case where residual_out is a alias of x, which makes torch.library
315
  # and torch.compile unhappy. Also allocate memory for out and residual_out if they are None
@@ -355,8 +376,11 @@ def _layer_norm_fwd(
355
 
356
  # [2025-04-28] torch.library.triton_op ignores the schema argument, but here we need the schema
357
  # since we're returning a tuple of tensors
358
- @triton_op(add_op_namespace_prefix("layer_norm_fwd_impl"), mutates_args={"out", "residual_out"},
359
- 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)")
 
 
 
360
  def _layer_norm_fwd_impl(
361
  x: Tensor,
362
  weight: Tensor,
@@ -372,7 +396,7 @@ def _layer_norm_fwd_impl(
372
  zero_centered_weight: bool = False,
373
  is_rms_norm: bool = False,
374
  return_dropout_mask: bool = False,
375
- residual_out: Optional[Tensor] = None
376
  ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
377
  M, N = x.shape
378
  assert x.stride(-1) == 1
@@ -407,7 +431,11 @@ def _layer_norm_fwd_impl(
407
  assert y1.stride(-1) == 1
408
  else:
409
  y1 = None
410
- mean = torch.empty((M,), dtype=torch.float32, device=x.device) if not is_rms_norm else None
 
 
 
 
411
  rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
412
  if dropout_p > 0.0:
413
  seeds = torch.randint(
@@ -475,7 +503,14 @@ def _layer_norm_fwd_impl(
475
 
476
  @triton.autotune(
477
  configs=triton_autotune_configs(),
478
- key=["N", "HAS_DRESIDUAL", "STORE_DRESIDUAL", "IS_RMS_NORM", "HAS_BIAS", "HAS_DROPOUT"],
 
 
 
 
 
 
 
479
  )
480
  # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
481
  # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
@@ -609,14 +644,18 @@ def _layer_norm_bwd_kernel(
609
  if HAS_DX1:
610
  if HAS_DROPOUT:
611
  keep_mask = (
612
- tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
 
613
  )
614
  dx1 = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
615
  else:
616
  dx1 = dx
617
  tl.store(DX1 + cols, dx1, mask=mask)
618
  if HAS_DROPOUT:
619
- keep_mask = tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
 
 
 
620
  dx = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
621
  if HAS_ROWSCALE:
622
  rowscale = tl.load(ROWSCALE + row).to(tl.float32)
@@ -699,11 +738,12 @@ def _layer_norm_bwd(
699
  return dx, dw, db, dresidual_in, dx1, dw1, db1, y
700
 
701
 
702
-
703
- @triton_op(add_op_namespace_prefix("layer_norm_bwd_impl"), mutates_args={},
704
- 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)",
705
- allow_decomposition=False, # Don't let torch.compile trace inside
706
- )
 
707
  def _layer_norm_bwd_impl(
708
  dy: Tensor,
709
  x: Tensor,
@@ -770,9 +810,15 @@ def _layer_norm_bwd_impl(
770
  else None
771
  )
772
  dx1 = torch.empty_like(dx) if (has_x1 and dropout_p > 0.0) else None
773
- y = torch.empty(M, N, dtype=dy.dtype, device=dy.device) if recompute_output else None
 
 
 
 
774
  if recompute_output:
775
- assert weight1 is None, "recompute_output is not supported with parallel LayerNorm"
 
 
776
 
777
  # Less than 64KB per feature: enqueue fused kernel
778
  MAX_FUSED_SIZE = 65536 // x.element_size()
@@ -849,7 +895,6 @@ def _layer_norm_bwd_impl(
849
 
850
 
851
  class LayerNormFn(torch.autograd.Function):
852
-
853
  @staticmethod
854
  def forward(
855
  ctx,
@@ -870,14 +915,16 @@ class LayerNormFn(torch.autograd.Function):
870
  return_dropout_mask=False,
871
  out_dtype=None,
872
  out=None,
873
- residual_out=None
874
  ):
875
  x_shape_og = x.shape
876
  # reshape input data into 2D tensor
877
  x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
878
  if residual is not None:
879
  assert residual.shape == x_shape_og
880
- residual = maybe_contiguous_lastdim(residual.reshape(-1, residual.shape[-1]))
 
 
881
  if x1 is not None:
882
  assert x1.shape == x_shape_og
883
  assert rowscale is None, "rowscale is not supported with parallel LayerNorm"
@@ -897,24 +944,26 @@ class LayerNormFn(torch.autograd.Function):
897
  out = out.reshape(-1, out.shape[-1])
898
  if residual_out is not None:
899
  residual_out = residual_out.reshape(-1, residual_out.shape[-1])
900
- y, y1, mean, rstd, residual_out, seeds, dropout_mask, dropout_mask1 = _layer_norm_fwd(
901
- x,
902
- weight,
903
- bias,
904
- eps,
905
- residual,
906
- x1,
907
- weight1,
908
- bias1,
909
- dropout_p=dropout_p,
910
- rowscale=rowscale,
911
- out_dtype=out_dtype,
912
- residual_dtype=residual_dtype,
913
- zero_centered_weight=zero_centered_weight,
914
- is_rms_norm=is_rms_norm,
915
- return_dropout_mask=return_dropout_mask,
916
- out=out,
917
- residual_out=residual_out,
 
 
918
  )
919
  ctx.save_for_backward(
920
  residual_out, weight, bias, weight1, bias1, rowscale, seeds, mean, rstd
@@ -930,9 +979,15 @@ class LayerNormFn(torch.autograd.Function):
930
  ctx.zero_centered_weight = zero_centered_weight
931
  y = y.reshape(x_shape_og)
932
  y1 = y1.reshape(x_shape_og) if y1 is not None else None
933
- residual_out = residual_out.reshape(x_shape_og) if residual_out is not None else None
934
- dropout_mask = dropout_mask.reshape(x_shape_og) if dropout_mask is not None else None
935
- dropout_mask1 = dropout_mask1.reshape(x_shape_og) if dropout_mask1 is not None else None
 
 
 
 
 
 
936
  if not return_dropout_mask:
937
  if weight1 is None:
938
  return y if not prenorm else (y, residual_out)
@@ -1030,7 +1085,7 @@ def layer_norm_fn(
1030
  return_dropout_mask=False,
1031
  out_dtype=None,
1032
  out=None,
1033
- residual_out=None
1034
  ):
1035
  return LayerNormFn.apply(
1036
  x,
@@ -1050,7 +1105,7 @@ def layer_norm_fn(
1050
  return_dropout_mask,
1051
  out_dtype,
1052
  out,
1053
- residual_out
1054
  )
1055
 
1056
 
@@ -1071,7 +1126,7 @@ def rms_norm_fn(
1071
  return_dropout_mask=False,
1072
  out_dtype=None,
1073
  out=None,
1074
- residual_out=None
1075
  ):
1076
  return LayerNormFn.apply(
1077
  x,
@@ -1091,14 +1146,20 @@ def rms_norm_fn(
1091
  return_dropout_mask,
1092
  out_dtype,
1093
  out,
1094
- residual_out
1095
  )
1096
 
1097
 
1098
  class RMSNorm(torch.nn.Module):
1099
-
1100
- def __init__(self, hidden_size, eps=1e-5, dropout_p=0.0, zero_centered_weight=False,
1101
- device=None, dtype=None):
 
 
 
 
 
 
1102
  factory_kwargs = {"device": device, "dtype": dtype}
1103
  super().__init__()
1104
  self.eps = eps
@@ -1132,7 +1193,6 @@ class RMSNorm(torch.nn.Module):
1132
 
1133
 
1134
  class LayerNormLinearFn(torch.autograd.Function):
1135
-
1136
  @staticmethod
1137
  @custom_fwd
1138
  def forward(
@@ -1153,7 +1213,9 @@ class LayerNormLinearFn(torch.autograd.Function):
1153
  x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
1154
  if residual is not None:
1155
  assert residual.shape == x_shape_og
1156
- residual = maybe_contiguous_lastdim(residual.reshape(-1, residual.shape[-1]))
 
 
1157
  norm_weight = norm_weight.contiguous()
1158
  norm_bias = maybe_contiguous(norm_bias)
1159
  residual_dtype = (
@@ -1167,17 +1229,23 @@ class LayerNormLinearFn(torch.autograd.Function):
1167
  norm_bias,
1168
  eps,
1169
  residual,
1170
- out_dtype=None if not torch.is_autocast_enabled() else torch.get_autocast_dtype("cuda"),
 
 
1171
  residual_dtype=residual_dtype,
1172
  is_rms_norm=is_rms_norm,
1173
  )
1174
  y = y.reshape(x_shape_og)
1175
- dtype = torch.get_autocast_dtype("cuda") if torch.is_autocast_enabled() else y.dtype
 
 
1176
  linear_weight = linear_weight.to(dtype)
1177
  linear_bias = linear_bias.to(dtype) if linear_bias is not None else None
1178
  out = F.linear(y.to(linear_weight.dtype), linear_weight, linear_bias)
1179
  # We don't store y, will be recomputed in the backward pass to save memory
1180
- ctx.save_for_backward(residual_out, norm_weight, norm_bias, linear_weight, mean, rstd)
 
 
1181
  ctx.x_shape_og = x_shape_og
1182
  ctx.eps = eps
1183
  ctx.is_rms_norm = is_rms_norm
@@ -1198,8 +1266,9 @@ class LayerNormLinearFn(torch.autograd.Function):
1198
  assert dy.shape == x.shape
1199
  if ctx.prenorm:
1200
  dresidual = args[0]
1201
- dresidual = maybe_contiguous_lastdim(dresidual.reshape(-1, dresidual.shape[-1]))
1202
- assert dresidual.shape == x.shape
 
1203
  else:
1204
  dresidual = None
1205
  dx, dnorm_weight, dnorm_bias, dresidual_in, _, _, _, y = _layer_norm_bwd(
 
7
  # The models we train have hidden dim up to 8k anyway (e.g. Llama 70B), so this is fine.
8
 
9
  import math
10
+ from typing import List, Optional
11
 
12
  import torch
13
  import torch.nn.functional as F
 
 
14
  import triton
15
  import triton.language as tl
16
+ from torch import Tensor
17
 
18
+ from ._ops import add_op_namespace_prefix
19
  from .utils.library import triton_op
20
+ from .utils.torch import custom_bwd, custom_fwd
21
 
22
 
23
  def maybe_contiguous_lastdim(x):
 
39
  warp_size = 32
40
  if torch.cuda.is_available():
41
  warp_size = getattr(torch.cuda.get_device_properties(torch.cuda.current_device()), "warp_size", 32)
42
+ warp_size = getattr(
43
+ torch.cuda.get_device_properties(torch.cuda.current_device()),
44
+ "warp_size",
45
+ 32,
46
+ )
47
  # Autotune for warp counts which are powers of 2 and do not exceed thread per block limit
48
+ return [
49
+ triton.Config({}, num_warps=warp_count)
50
+ for warp_count in [1, 2, 4, 8, 16, 32]
51
+ if warp_count * warp_size <= max_threads_per_block
52
+ ]
53
  # return [triton.Config({}, num_warps=8)]
54
 
55
 
 
101
  x = x + x1
102
  if residual is not None:
103
  x = (x + residual).to(x.dtype)
104
+ out = F.layer_norm(
105
+ x.to(weight.dtype), x.shape[-1:], weight=weight, bias=bias, eps=eps
106
+ ).to(dtype)
107
  if weight1 is None:
108
  return out if not prenorm else (out, x)
109
  else:
 
162
  if residual is not None:
163
  x = (x + residual).to(x.dtype)
164
  rstd = 1 / torch.sqrt((x.square()).mean(dim=-1, keepdim=True) + eps)
165
+ out = ((x * rstd * weight) + bias if bias is not None else (x * rstd * weight)).to(
166
+ dtype
167
+ )
168
  if weight1 is None:
169
  return out if not prenorm else (out, x)
170
  else:
171
+ out1 = (
172
+ (x * rstd * weight1) + bias1 if bias1 is not None else (x * rstd * weight1)
173
+ ).to(dtype)
174
  return (out, out1) if not prenorm else (out, out1, x)
175
 
176
 
177
  @triton.autotune(
178
  configs=triton_autotune_configs(),
179
+ key=[
180
+ "N",
181
+ "HAS_RESIDUAL",
182
+ "STORE_RESIDUAL_OUT",
183
+ "IS_RMS_NORM",
184
+ "HAS_BIAS",
185
+ "HAS_X1",
186
+ "HAS_W1",
187
+ "HAS_B1",
188
+ ],
189
  )
190
  # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
191
  # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
 
255
  if HAS_DROPOUT:
256
  # Compute dropout mask
257
  # 7 rounds is good enough, and reduces register pressure
258
+ keep_mask = (
259
+ tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7) > dropout_p
260
+ )
261
  x = tl.where(keep_mask, x / (1.0 - dropout_p), 0.0)
262
  if STORE_DROPOUT_MASK:
263
  tl.store(DROPOUT_MASK + row * N + cols, keep_mask, mask=cols < N)
 
270
  # Compute dropout mask
271
  # 7 rounds is good enough, and reduces register pressure
272
  keep_mask = (
273
+ tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7)
274
+ > dropout_p
275
  )
276
  x1 = tl.where(keep_mask, x1 / (1.0 - dropout_p), 0.0)
277
  if STORE_DROPOUT_MASK:
 
330
  is_rms_norm: bool = False,
331
  return_dropout_mask: bool = False,
332
  out: Optional[Tensor] = None,
333
+ residual_out: Optional[Tensor] = None,
334
  ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
335
  # Need to wrap to handle the case where residual_out is a alias of x, which makes torch.library
336
  # and torch.compile unhappy. Also allocate memory for out and residual_out if they are None
 
376
 
377
  # [2025-04-28] torch.library.triton_op ignores the schema argument, but here we need the schema
378
  # since we're returning a tuple of tensors
379
+ @triton_op(
380
+ add_op_namespace_prefix("layer_norm_fwd_impl"),
381
+ mutates_args={"out", "residual_out"},
382
+ 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)",
383
+ )
384
  def _layer_norm_fwd_impl(
385
  x: Tensor,
386
  weight: Tensor,
 
396
  zero_centered_weight: bool = False,
397
  is_rms_norm: bool = False,
398
  return_dropout_mask: bool = False,
399
+ residual_out: Optional[Tensor] = None,
400
  ) -> (Tensor, Tensor, Tensor, Tensor, Tensor, Tensor):
401
  M, N = x.shape
402
  assert x.stride(-1) == 1
 
431
  assert y1.stride(-1) == 1
432
  else:
433
  y1 = None
434
+ mean = (
435
+ torch.empty((M,), dtype=torch.float32, device=x.device)
436
+ if not is_rms_norm
437
+ else None
438
+ )
439
  rstd = torch.empty((M,), dtype=torch.float32, device=x.device)
440
  if dropout_p > 0.0:
441
  seeds = torch.randint(
 
503
 
504
  @triton.autotune(
505
  configs=triton_autotune_configs(),
506
+ key=[
507
+ "N",
508
+ "HAS_DRESIDUAL",
509
+ "STORE_DRESIDUAL",
510
+ "IS_RMS_NORM",
511
+ "HAS_BIAS",
512
+ "HAS_DROPOUT",
513
+ ],
514
  )
515
  # torch compile doesn't like triton.heuristics, so we set these manually when calling the kernel
516
  # @triton.heuristics({"HAS_BIAS": lambda args: args["B"] is not None})
 
644
  if HAS_DX1:
645
  if HAS_DROPOUT:
646
  keep_mask = (
647
+ tl.rand(tl.load(SEEDS + M + row).to(tl.uint32), cols, n_rounds=7)
648
+ > dropout_p
649
  )
650
  dx1 = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
651
  else:
652
  dx1 = dx
653
  tl.store(DX1 + cols, dx1, mask=mask)
654
  if HAS_DROPOUT:
655
+ keep_mask = (
656
+ tl.rand(tl.load(SEEDS + row).to(tl.uint32), cols, n_rounds=7)
657
+ > dropout_p
658
+ )
659
  dx = tl.where(keep_mask, dx / (1.0 - dropout_p), 0.0)
660
  if HAS_ROWSCALE:
661
  rowscale = tl.load(ROWSCALE + row).to(tl.float32)
 
738
  return dx, dw, db, dresidual_in, dx1, dw1, db1, y
739
 
740
 
741
+ @triton_op(
742
+ add_op_namespace_prefix("layer_norm_bwd_impl"),
743
+ mutates_args={},
744
+ 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)",
745
+ allow_decomposition=False, # Don't let torch.compile trace inside
746
+ )
747
  def _layer_norm_bwd_impl(
748
  dy: Tensor,
749
  x: Tensor,
 
810
  else None
811
  )
812
  dx1 = torch.empty_like(dx) if (has_x1 and dropout_p > 0.0) else None
813
+ y = (
814
+ torch.empty(M, N, dtype=dy.dtype, device=dy.device)
815
+ if recompute_output
816
+ else None
817
+ )
818
  if recompute_output:
819
+ assert weight1 is None, (
820
+ "recompute_output is not supported with parallel LayerNorm"
821
+ )
822
 
823
  # Less than 64KB per feature: enqueue fused kernel
824
  MAX_FUSED_SIZE = 65536 // x.element_size()
 
895
 
896
 
897
  class LayerNormFn(torch.autograd.Function):
 
898
  @staticmethod
899
  def forward(
900
  ctx,
 
915
  return_dropout_mask=False,
916
  out_dtype=None,
917
  out=None,
918
+ residual_out=None,
919
  ):
920
  x_shape_og = x.shape
921
  # reshape input data into 2D tensor
922
  x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
923
  if residual is not None:
924
  assert residual.shape == x_shape_og
925
+ residual = maybe_contiguous_lastdim(
926
+ residual.reshape(-1, residual.shape[-1])
927
+ )
928
  if x1 is not None:
929
  assert x1.shape == x_shape_og
930
  assert rowscale is None, "rowscale is not supported with parallel LayerNorm"
 
944
  out = out.reshape(-1, out.shape[-1])
945
  if residual_out is not None:
946
  residual_out = residual_out.reshape(-1, residual_out.shape[-1])
947
+ y, y1, mean, rstd, residual_out, seeds, dropout_mask, dropout_mask1 = (
948
+ _layer_norm_fwd(
949
+ x,
950
+ weight,
951
+ bias,
952
+ eps,
953
+ residual,
954
+ x1,
955
+ weight1,
956
+ bias1,
957
+ dropout_p=dropout_p,
958
+ rowscale=rowscale,
959
+ out_dtype=out_dtype,
960
+ residual_dtype=residual_dtype,
961
+ zero_centered_weight=zero_centered_weight,
962
+ is_rms_norm=is_rms_norm,
963
+ return_dropout_mask=return_dropout_mask,
964
+ out=out,
965
+ residual_out=residual_out,
966
+ )
967
  )
968
  ctx.save_for_backward(
969
  residual_out, weight, bias, weight1, bias1, rowscale, seeds, mean, rstd
 
979
  ctx.zero_centered_weight = zero_centered_weight
980
  y = y.reshape(x_shape_og)
981
  y1 = y1.reshape(x_shape_og) if y1 is not None else None
982
+ residual_out = (
983
+ residual_out.reshape(x_shape_og) if residual_out is not None else None
984
+ )
985
+ dropout_mask = (
986
+ dropout_mask.reshape(x_shape_og) if dropout_mask is not None else None
987
+ )
988
+ dropout_mask1 = (
989
+ dropout_mask1.reshape(x_shape_og) if dropout_mask1 is not None else None
990
+ )
991
  if not return_dropout_mask:
992
  if weight1 is None:
993
  return y if not prenorm else (y, residual_out)
 
1085
  return_dropout_mask=False,
1086
  out_dtype=None,
1087
  out=None,
1088
+ residual_out=None,
1089
  ):
1090
  return LayerNormFn.apply(
1091
  x,
 
1105
  return_dropout_mask,
1106
  out_dtype,
1107
  out,
1108
+ residual_out,
1109
  )
1110
 
1111
 
 
1126
  return_dropout_mask=False,
1127
  out_dtype=None,
1128
  out=None,
1129
+ residual_out=None,
1130
  ):
1131
  return LayerNormFn.apply(
1132
  x,
 
1146
  return_dropout_mask,
1147
  out_dtype,
1148
  out,
1149
+ residual_out,
1150
  )
1151
 
1152
 
1153
  class RMSNorm(torch.nn.Module):
1154
+ def __init__(
1155
+ self,
1156
+ hidden_size,
1157
+ eps=1e-5,
1158
+ dropout_p=0.0,
1159
+ zero_centered_weight=False,
1160
+ device=None,
1161
+ dtype=None,
1162
+ ):
1163
  factory_kwargs = {"device": device, "dtype": dtype}
1164
  super().__init__()
1165
  self.eps = eps
 
1193
 
1194
 
1195
  class LayerNormLinearFn(torch.autograd.Function):
 
1196
  @staticmethod
1197
  @custom_fwd
1198
  def forward(
 
1213
  x = maybe_contiguous_lastdim(x.reshape(-1, x.shape[-1]))
1214
  if residual is not None:
1215
  assert residual.shape == x_shape_og
1216
+ residual = maybe_contiguous_lastdim(
1217
+ residual.reshape(-1, residual.shape[-1])
1218
+ )
1219
  norm_weight = norm_weight.contiguous()
1220
  norm_bias = maybe_contiguous(norm_bias)
1221
  residual_dtype = (
 
1229
  norm_bias,
1230
  eps,
1231
  residual,
1232
+ out_dtype=None
1233
+ if not torch.is_autocast_enabled()
1234
+ else torch.get_autocast_dtype("cuda"),
1235
  residual_dtype=residual_dtype,
1236
  is_rms_norm=is_rms_norm,
1237
  )
1238
  y = y.reshape(x_shape_og)
1239
+ dtype = (
1240
+ torch.get_autocast_dtype("cuda") if torch.is_autocast_enabled() else y.dtype
1241
+ )
1242
  linear_weight = linear_weight.to(dtype)
1243
  linear_bias = linear_bias.to(dtype) if linear_bias is not None else None
1244
  out = F.linear(y.to(linear_weight.dtype), linear_weight, linear_bias)
1245
  # We don't store y, will be recomputed in the backward pass to save memory
1246
+ ctx.save_for_backward(
1247
+ residual_out, norm_weight, norm_bias, linear_weight, mean, rstd
1248
+ )
1249
  ctx.x_shape_og = x_shape_og
1250
  ctx.eps = eps
1251
  ctx.is_rms_norm = is_rms_norm
 
1266
  assert dy.shape == x.shape
1267
  if ctx.prenorm:
1268
  dresidual = args[0]
1269
+ dresidual = maybe_contiguous_lastdim(
1270
+ dresidual.reshape(-1, dresidual.shape[-1])
1271
+ )
1272
  else:
1273
  dresidual = None
1274
  dx, dnorm_weight, dnorm_bias, dresidual_in, _, _, _, y = _layer_norm_bwd(
build/torch-xpu/metadata.json CHANGED
@@ -1,6 +1,6 @@
1
  {
2
  "name": "flash-attn-ops",
3
- "id": "_flash_attn_ops_xpu_2e4b976",
4
  "version": 1,
5
  "license": "BSD-3-Clause",
6
  "python-depends": [],
@@ -11,17 +11,27 @@
11
  "algorithm": "sha256",
12
  "files": {
13
  "__init__.py": "CN2UymE7gbYnzhbAm92qHBHVC/as55EjJsW+EQt5eaM=",
14
- "_ops.py": "p2iM6XDGC6V+jZxNrRc7AwioU+OaCDlgkyHbF++T5Vg=",
15
- "_ops_compat.py": "xs6BjkzWExaqXt9xniN5UNYqTp0kpEehRyOyx16+a5o=",
16
  "cross_entropy.py": "y56nmuOtpSNxM7sDZcUSgXKIsIEayfuvrU8bAVog/nw=",
17
  "flash_attn_ops/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
18
  "k_activations.py": "+Z3vIyO4JkqBMipKsPvhzmxljtBdIhJCsl/M+/ESqBo=",
19
- "layer_norm.py": "BTwZtyGUbjmWsox0rbMomXsc/UaZJAs+4iexOJ2PfIU=",
20
  "losses.py": "YhTM9q3Qouq8QoFy1hryAXFC2A6scVtnReoNa+l99YQ=",
21
  "rotary.py": "0gN+YjhUxwL7B1eE/kKS1Kida5rBrjPKs9Z4j4+U8kA=",
22
  "utils/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
23
  "utils/library.py": "nKTJmAHNSqc30g95/NPDtvqzP4GWgYN0BNh015AFzYM=",
24
  "utils/torch.py": "7mVlWFrCGE7XuCLO9oHa5Z8c4HPLYnE0ojYOJHh8JiQ="
25
  }
 
 
 
 
 
 
 
 
 
 
 
26
  }
27
  }
 
1
  {
2
  "name": "flash-attn-ops",
3
+ "id": "_flash_attn_ops_xpu_8a730d9",
4
  "version": 1,
5
  "license": "BSD-3-Clause",
6
  "python-depends": [],
 
11
  "algorithm": "sha256",
12
  "files": {
13
  "__init__.py": "CN2UymE7gbYnzhbAm92qHBHVC/as55EjJsW+EQt5eaM=",
14
+ "_ops.py": "qQNowLEFyFPzxnfTsNt9k+zqYB65oQJozdV6LFvAvNQ=",
 
15
  "cross_entropy.py": "y56nmuOtpSNxM7sDZcUSgXKIsIEayfuvrU8bAVog/nw=",
16
  "flash_attn_ops/__init__.py": "DFYPlrhXwYjEqCl/8n0SmWGZV8NFml5DPhMjKfv98GY=",
17
  "k_activations.py": "+Z3vIyO4JkqBMipKsPvhzmxljtBdIhJCsl/M+/ESqBo=",
18
+ "layer_norm.py": "SVUTedaQF7zdKwg4bNnr8NPWFSUPB+88kSLWyPkykh0=",
19
  "losses.py": "YhTM9q3Qouq8QoFy1hryAXFC2A6scVtnReoNa+l99YQ=",
20
  "rotary.py": "0gN+YjhUxwL7B1eE/kKS1Kida5rBrjPKs9Z4j4+U8kA=",
21
  "utils/__init__.py": "47DEQpj8HBSa+/TImW+5JCeuQeRkm5NMpJWZG3hSuFU=",
22
  "utils/library.py": "nKTJmAHNSqc30g95/NPDtvqzP4GWgYN0BNh015AFzYM=",
23
  "utils/torch.py": "7mVlWFrCGE7XuCLO9oHa5Z8c4HPLYnE0ojYOJHh8JiQ="
24
  }
25
+ },
26
+ "provenance": {
27
+ "kernel-builder": {
28
+ "version": "0.17.0-dev0",
29
+ "sha": "81580bb92577f2f7228661ca2a221fc052375709",
30
+ "dirty": false
31
+ },
32
+ "kernel": {
33
+ "sha": "8a730d96c37560ccf1d3e09f7bbccdc886818f33",
34
+ "dirty": false
35
+ }
36
  }
37
  }
build/torch-xpu/metadata.json.sigstore CHANGED
@@ -1 +1 @@
1
- {"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json","verificationMaterial":{"certificate":{"rawBytes":"MIIHSzCCBtKgAwIBAgIUF8gMxpXCQneYw620L9/qmEPn348wCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwNzAzMTIwNzEyWhcNMjYwNzAzMTIxNzEyWjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAENdW72gOKARAvW9ciCNkfeEWQLNjvhtbhKKVMYUJEwBOz9xkKy5ih8nK54uNO4JOhnnilThdpi4gWyQy9An7zrKOCBfEwggXtMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUB49iHn6qMC4ep475g778ZZoaog0wHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoMmU0Yjk3NjIxZTkzNzk3NDEzMmY1NDNkYjI1MTEyNjgwNmY5NGJiZTAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDJlNGI5NzYyMWU5Mzc5NzQxMzJmNTQzZGIyNTExMjY4MDZmOTRiYmUwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMjg2NTg3NTI1MjgvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBiwYKKwYBBAHWeQIEAgR9BHsAeQB3AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABnyfgMo8AAAQDAEgwRgIhAMCtyvaiH/9vCbWotbwTvUhke9tJr9M8S5S7zJ232UatAiEAxK8BxOXUhKN5g5FgBoQD9BvKII41Wn1S3S9cVQp0DHIwCgYIKoZIzj0EAwMDZwAwZAIwB3Hdfv/uX2SOdgFV8PqkENmZ0UFeCkr8QhkcOVMPV7RDvk39qsT3zLPYleIK+G7nAjAcs5CGZlTUV3D2bPyyS6LKeGe80M+Qg9vw3IIjnkM60Caj8p36D1wAkL9dITaHb2k="},"tlogEntries":[{"logIndex":"2060168484","logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="},"kindVersion":{"kind":"hashedrekord","version":"0.0.1"},"integratedTime":"1783080433","inclusionPromise":{"signedEntryTimestamp":"MEUCIDMuRXbeQSfMrRmWGx0jZyGr6RnxDwUVErjUrVc1xNARAiEAmA2L+7PTjJXnnIom1w9gXuWJ72z4b++3dzlhqYGboIc="},"inclusionProof":{"logIndex":"1938264222","rootHash":"XJis1D7kR+9qFOiIvVh6o6xO2882B0g6JNwolzSAal8=","treeSize":"1938264253","hashes":["0rnr0+f4ipe8UokecBe7O+X4mesnXw5CQ3H8Atii/Ts=","W7KZMwpfsKmqz5BbnUZkrFD95AicOWA6dEX0vk2YR4c=","e82RbLiB6NiR7Qvw4lGabr7gu+9XX3oX6j8XOGn+xV4=","/JUlk9idNhwWXmOAHjOEnNBdMOzm8HsbTDjftqtLF5I=","yaj1mvlBRD39f6VKsE6/Sk+hUn10D6W6AQI1nAwhftk=","+5r7J2Rca3nsJENDrEDGluJEwjRd5zybZCMxkx02XQM=","pGx12HyM81J7B68eP7kLLZuRUSPNyuzw1KpBPHvDjxE=","y9pMOMTBiXS0/nNsWl6km9Rbt47KgljMyLqu/0Rvm1M=","kdmXwq0Vx67sWzz9xJdfxjd6WXw6nNtrpUOt2BYrQPg=","4/a1MJsWpAoolCIszNdX3KIuhcFLrSbhvmOzjO83muI=","MjeWaEt50oS4ciA2SzTy+1gTRd7OQ+MIQadc4CXFv3I=","JYl6P0ZznUGRvcZdNHZqKkY3KNFNz/UowViPyJQT0mg=","RIZjwQJv61lwM5JOP8iNT/RxKPr522/etVqxVBWn7KA=","b69YlhA5qlDwN70us4bq9ysnikE1ulA7EwKBTb+atY0=","LM7F5iwbk8jG29X4FwpkCN6BqDCfuNQTlz5CyQWt87s=","+/VZ56MsIPxMiyLAodzKXo5TEWdQp36z89qLhpzloAo=","daxmZaajRpZV+JxHiOYZhJBiSKN5ucqjh2WnGbHhirw=","DOCeoSMovIvLExkhIvisow9AuNXgeWs4ECkyR6EcqYU="],"checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n1938264253\nXJis1D7kR+9qFOiIvVh6o6xO2882B0g6JNwolzSAal8=\n\n— rekor.sigstore.dev wNI9ajBFAiEAjUCRuYyZp5ol+iq1JyAkcOH8oQfUeYskcgWPMC7OeAcCIEKmYjNRbUrTXzEDuvytOfidjE51hHrmoMy3iNiJ9kJ2\n"}},"canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJiZmQ0OGJjMmFhZmMzNWRhMmU2YTVmNGQ0M2Q3M2YzM2EyMzM2MTQyN2FkN2QyNjA2N2UyMTNiNjhiN2FhYTkwIn19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FWUNJUURZanV3ODhYT1JPT0NMRG5sM2V0bThqTFhHUW1jbTdFSnJLajAyYjhwTXV3SWhBT1hHOUFkb1dZU3RGYVRpUzdRSDB4MStjbDlhQWxtRG9WaEdHbVZ5YTZrMSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRla05EUW5STFowRjNTVUpCWjBsVlJqaG5UWGh3V0VOUmJtVlpkell5TUV3NUwzRnRSVkJ1TXpRNGQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDU2UVhwTlZFbDNUbnBGZVZkb1kwNU5hbGwzVG5wQmVrMVVTWGhPZWtWNVYycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZPWkZjM01tZFBTMEZTUVhaWE9XTnBRMDVyWm1WRlYxRk1UbXAyYUhSaWFFdExWazBLV1ZWS1JYZENUM281ZUd0TGVUVnBhRGh1U3pVMGRVNVBORXBQYUc1dWFXeFVhR1J3YVRSblYzbFJlVGxCYmpkNmNrdFBRMEptUlhkbloxaDBUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZDTkRscENraHVObkZOUXpSbGNEUTNOV2MzTnpoYVdtOWhiMmN3ZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOU5iVlV3V1dwck0wNXFTWGhhVkd0NlRucHJNMDVFUlhwTmJWa3hDazVFVG10WmFra3hUVlJGZVU1cVozZE9iVmsxVGtkS2FWcFVRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDF0VlRCWmFtc3pUbXBKZUZwVWEzcE9lbXN6VGtSRmVrMXRXVEZPUkU1cldXcEpNVTFVUlhrS1RtcG5kMDV0V1RWT1IwcHBXbFJCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDF0VlRCWmFtc3pUbXBKZUZwVWEzcE9lbXN6VGtSRmVrMXRXVEVLVGtST2ExbHFTVEZOVkVWNVRtcG5kMDV0V1RWT1IwcHBXbFJCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUktiRTVIU1RVS1RucFplVTFYVlRWTmVtTTFUbnBSZUUxNlNtMU9WRkY2V2tkSmVVNVVSWGhOYWxrMFRVUmFiVTlVVW1sWmJWVjNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOYW1jeVRsUm5NMDVVU1RGTmFtZDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwZDFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT1VKSWMwRUtaVkZDTTBGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtNTVabWROYnpoQlFVRlJSQXBCUldkM1VtZEphRUZOUTNSNWRtRnBTQzg1ZGtOaVYyOTBZbmRVZGxWb2EyVTVkRXB5T1UwNFV6VlROM3BLTWpNeVZXRjBRV2xGUVhoTE9FSjRUMWhWQ21oTFRqVm5OVVpuUW05UlJEbENka3RKU1RReFYyNHhVek5UT1dOV1VYQXdSRWhKZDBObldVbExiMXBKZW1vd1JVRjNUVVJhZDBGM1drRkpkMEl6U0dRS1puWXZkVmd5VTA5a1owWldPRkJ4YTBWT2JWb3dWVVpsUTJ0eU9GRm9hMk5QVmsxUVZqZFNSSFpyTXpseGMxUXpla3hRV1d4bFNVc3JSemR1UVdwQll3cHpOVU5IV214VVZWWXpSREppVUhsNVV6Wk1TMlZIWlRnd1RTdFJaemwyZHpOSlNXcHVhMDAyTUVOaGFqaHdNelpFTVhkQmEwdzVaRWxVWVVoaU1tczlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}],"timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbcGCyqGSIb3DQEJEAEEoIGnBIGkMIGhAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgw51+ID7l4Ye3B4BTtSYIbJcGJlyn0z2q9vG5bpbnFnECFH3/8ZKUxZU4kuXxHIKLid+nHHrPGA8yMDI2MDcwMzEyMDcxMlowAwIBAaAypDAwLjEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MRUwEwYDVQQDEwxzaWdzdG9yZS10c2GgADGCAdswggHXAgEBMFEwOTEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MSAwHgYDVQQDExdzaWdzdG9yZS10c2Etc2VsZnNpZ25lZAIUOhNULwyQYe68wUMvy4qOiyojiwwwCwYJYIZIAWUDBAIBoIH8MBoGCSqGSIb3DQEJAzENBgsqhkiG9w0BCRABBDAcBgkqhkiG9w0BCQUxDxcNMjYwNzAzMTIwNzEyWjAvBgkqhkiG9w0BCQQxIgQgI0MlPBBdHm5i0ufPV3GtTjhbzXn3nQJovwyVAYTYongwgY4GCyqGSIb3DQEJEAIvMX8wfTB7MHkEIIX5J7wHq2LKw7RDVsEO/IGyxog/2nq55thw2dE6zQW3MFUwPaQ7MDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAoGCCqGSM49BAMCBGcwZQIwTOA+nOVYQyQ+tBQagjVXCEQANh5c7LJ1cahi3edm/jyQWxu2tIaMj1J8Ac+kARA3AjEAvtcanpIFrMVFms7U8YE9y3xGXo1OoT9ix3xsO3OOOgSXrW9Jb0npKE5NjbBdgImY"}]}},"messageSignature":{"messageDigest":{"algorithm":"SHA2_256","digest":"v9SLwqr8Ndoual9NQ9c/M6IzYUJ619JgZ+ITtot6qpA="},"signature":"MEYCIQDYjuw88XOROOCLDnl3etm8jLXGQmcm7EJrKj02b8pMuwIhAOXG9AdoWYStFaTiS7QH0x1+cl9aAlmDoVhGGmVya6k1"}}
 
1
+ {"mediaType":"application/vnd.dev.sigstore.bundle.v0.3+json","verificationMaterial":{"certificate":{"rawBytes":"MIIHSjCCBtGgAwIBAgIUEtkm7K/Pr9nzTXgrEfQdPMpsXCMwCgYIKoZIzj0EAwMwNzEVMBMGA1UEChMMc2lnc3RvcmUuZGV2MR4wHAYDVQQDExVzaWdzdG9yZS1pbnRlcm1lZGlhdGUwHhcNMjYwNzA4MTEyMTQ0WhcNMjYwNzA4MTEzMTQ0WjAAMFkwEwYHKoZIzj0CAQYIKoZIzj0DAQcDQgAEetuC8sjowRxXjIHUwVdgbjFqnEaeKxvIC8IG03knO5mDgMdBbp3YPXiQVKi9ERYHafd4ryooJpgOTBp+1VMv76OCBfAwggXsMA4GA1UdDwEB/wQEAwIHgDATBgNVHSUEDDAKBggrBgEFBQcDAzAdBgNVHQ4EFgQUbWOOmtt532eL9UmBTAOwy7WSfsMwHwYDVR0jBBgwFoAU39Ppz1YkEZb5qNjpKFWixi4YZD8wawYDVR0RAQH/BGEwX4ZdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDkGCisGAQQBg78wAQEEK2h0dHBzOi8vdG9rZW4uYWN0aW9ucy5naXRodWJ1c2VyY29udGVudC5jb20wHwYKKwYBBAGDvzABAgQRd29ya2Zsb3dfZGlzcGF0Y2gwNgYKKwYBBAGDvzABAwQoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzATBgorBgEEAYO/MAEEBAVCdWlsZDArBgorBgEEAYO/MAEFBB1odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eTAdBgorBgEEAYO/MAEGBA9yZWZzL2hlYWRzL21haW4wOwYKKwYBBAGDvzABCAQtDCtodHRwczovL3Rva2VuLmFjdGlvbnMuZ2l0aHVidXNlcmNvbnRlbnQuY29tMG0GCisGAQQBg78wAQkEXwxdaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5Ly5naXRodWIvd29ya2Zsb3dzL2J1aWxkLnlhbWxAcmVmcy9oZWFkcy9tYWluMDgGCisGAQQBg78wAQoEKgwoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzAbBgorBgEEAYO/MAELBA0MC3NlbGYtaG9zdGVkMEAGCisGAQQBg78wAQwEMgwwaHR0cHM6Ly9naXRodWIuY29tL2h1Z2dpbmdmYWNlL2tlcm5lbHMtY29tbXVuaXR5MDgGCisGAQQBg78wAQ0EKgwoOGE3MzBkOTZjMzc1NjBjY2YxZDNlMDlmN2JiY2NkYzg4NjgxOGYzMzAfBgorBgEEAYO/MAEOBBEMD3JlZnMvaGVhZHMvbWFpbjAaBgorBgEEAYO/MAEPBAwMCjEwNzE0NzU1MjkwLgYKKwYBBAGDvzABEAQgDB5odHRwczovL2dpdGh1Yi5jb20vaHVnZ2luZ2ZhY2UwGAYKKwYBBAGDvzABEQQKDAgyNTcyMDc0MzBtBgorBgEEAYO/MAESBF8MXWh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS8uZ2l0aHViL3dvcmtmbG93cy9idWlsZC55YW1sQHJlZnMvaGVhZHMvbWFpbjA4BgorBgEEAYO/MAETBCoMKDhhNzMwZDk2YzM3NTYwY2NmMWQzZTA5ZjdiYmNjZGM4ODY4MThmMzMwIQYKKwYBBAGDvzABFAQTDBF3b3JrZmxvd19kaXNwYXRjaDBkBgorBgEEAYO/MAEVBFYMVGh0dHBzOi8vZ2l0aHViLmNvbS9odWdnaW5nZmFjZS9rZXJuZWxzLWNvbW11bml0eS9hY3Rpb25zL3J1bnMvMjg5MzEyMDI0MzIvYXR0ZW1wdHMvMTAWBgorBgEEAYO/MAEWBAgMBnB1YmxpYzBGBgorBgEEAYO/MAEYBDgMNnJlcG86aHVnZ2luZ2ZhY2Uva2VybmVscy1jb21tdW5pdHk6cmVmOnJlZnMvaGVhZHMvbWFpbjCBigYKKwYBBAHWeQIEAgR8BHoAeAB2AN09MGrGxxEyYxkeHJlnNwKiSl643jyt/4eKcoAvKe6OAAABn0F2YR0AAAQDAEcwRQIgOJWSEuOyqL9n1RdMb3ARdqTJiOM3HaL0mIdaw1kJaskCIQCUuAp5nnEoOP3SWjFptQHRv3ruld6/KbLn9DtEKTDlCDAKBggqhkjOPQQDAwNnADBkAjBnS2Og6FZIQgGC7akGbABWrM2AriaUyUeA5k0cXdDtT2E02akrfahJCWXDbWJDKNwCMG9qUc2vBTHckCM0HwtutHq5GC9cQFBp26kz4UNzQyITdZPjZO29jUcw0RAIZx0/aQ=="},"tlogEntries":[{"logIndex":"2115129272","logId":{"keyId":"wNI9atQGlz+VWfO6LRygH4QUfY/8W4RFwiT5i5WRgB0="},"kindVersion":{"kind":"hashedrekord","version":"0.0.1"},"integratedTime":"1783509705","inclusionPromise":{"signedEntryTimestamp":"MEUCIH0ff510Evwca1u+UbwixmJ05HHv05MRu2mWjmWgoKx8AiEA7NCHpn7XlpJmWyxQ2l9PgDPrx/KxzgQRmzZm1pi/7co="},"inclusionProof":{"logIndex":"1993225010","rootHash":"3El3Lpy1oF3IYJ5Q+qyhZUxqDJqBzgrQpiLNQIitUkY=","treeSize":"1993225022","hashes":["xLUd2uZrFZnty9Z1HY2boG2lh+BB9B/UxRgRZgjFbuQ=","GSRgOWkhoIa6E3L/Xvm6sqXgl9omA9UrQLBwOH+9zbo=","wNw4VX/0O2nlFKgbRG0nQVgXHg2ywyCzflDmaiaglzI=","VnpjCIvL7brc6PK1dn48+tHR50n7RqdoDWrj5RzD1jE=","+7duWj64Qe1Zkx+yM9ISn5JVVI+HmTMcw3XpTLj8PPs=","446U/YKahjRVCpv9Upm9MAcm/hyhXbMT0TJora3feN8=","Q8PApMjMRUNJ7xrZy4f5UvQCyHXEOn4ZZ4p5sDiDc2g=","bi8fe0DDdrNAeoQsjKJr5a1b59xkRDANI6WC/eTjVno=","QllT6LSG/atuuaIF8rT3cyFz3Vv6oUwojQFAhFEwZl4=","p2o5Ruq9tV5P95j9zl6fx3XmOn49Uz6ZdywQA2WgVss=","VgOceJvrniBf3NxYSnd8hKhKiA3kzGYv9aw1wOrHMq8=","KEoaBedJtDHZJJ0hfsWskdmVy2dzbHOA8Z/QRgDoCAw=","Qt1jiKhOtf6yp73wykIuRUIfatVXkF3q7xBJ6DqwXnk=","frI9BMUdqa/KPLUCZ9yfDzZ2Z6tglf7SvThMJqEFBKU=","Bcvbbu+tDL+zr1H2q26c2ATmVu7v/f7JIFweUu7qDK8=","WajL3ZhY6ZX9meV3VTH9TkrbM31OF4o6+7foL0Uf6PI=","cyKMKBvHkQtwb2xR5svh+YppANOVdn7X8t4DuKMxICA=","+/VZ56MsIPxMiyLAodzKXo5TEWdQp36z89qLhpzloAo=","daxmZaajRpZV+JxHiOYZhJBiSKN5ucqjh2WnGbHhirw=","DOCeoSMovIvLExkhIvisow9AuNXgeWs4ECkyR6EcqYU="],"checkpoint":{"envelope":"rekor.sigstore.dev - 1193050959916656506\n1993225022\n3El3Lpy1oF3IYJ5Q+qyhZUxqDJqBzgrQpiLNQIitUkY=\n\n— rekor.sigstore.dev wNI9ajBEAiBL1aZWUnFZOHlpGow0mrf2+tBYoQvDfJWhDr1R59ADNwIgQU+XCes6bT3zIBuYPLQEnw42YWy5FbuduvCEQX7qbe0=\n"}},"canonicalizedBody":"eyJhcGlWZXJzaW9uIjoiMC4wLjEiLCJraW5kIjoiaGFzaGVkcmVrb3JkIiwic3BlYyI6eyJkYXRhIjp7Imhhc2giOnsiYWxnb3JpdGhtIjoic2hhMjU2IiwidmFsdWUiOiJjYTM3ZGYxN2RlNmQzYWRhNTMyMDEyNDE1YTZiNGM4OTQ1MmZhZjc0YWViY2YzM2Q2NzlhOGQ4M2Y1Yjg0ZjY1In19LCJzaWduYXR1cmUiOnsiY29udGVudCI6Ik1FUUNJQWxkQjFmTzJrWFQ1aEwwQWQvYjhkQmgxQWw4MHAzRVk3YUNhYmt0N3pXSUFpQjZCdnBRUnZnUG1GZ2MrZzV2WDVBSzZmSmVCUmxQaU55NFdlR1o4N2hVZGc9PSIsInB1YmxpY0tleSI6eyJjb250ZW50IjoiTFMwdExTMUNSVWRKVGlCRFJWSlVTVVpKUTBGVVJTMHRMUzB0Q2sxSlNVaFRha05EUW5SSFowRjNTVUpCWjBsVlJYUnJiVGRMTDFCeU9XNTZWRmhuY2tWbVVXUlFUWEJ6V0VOTmQwTm5XVWxMYjFwSmVtb3dSVUYzVFhjS1RucEZWazFDVFVkQk1WVkZRMmhOVFdNeWJHNWpNMUoyWTIxVmRWcEhWakpOVWpSM1NFRlpSRlpSVVVSRmVGWjZZVmRrZW1SSE9YbGFVekZ3WW01U2JBcGpiVEZzV2tkc2FHUkhWWGRJYUdOT1RXcFpkMDU2UVRSTlZFVjVUVlJSTUZkb1kwNU5hbGwzVG5wQk5FMVVSWHBOVkZFd1YycEJRVTFHYTNkRmQxbElDa3R2V2tsNmFqQkRRVkZaU1V0dldrbDZhakJFUVZGalJGRm5RVVZsZEhWRE9ITnFiM2RTZUZocVNVaFZkMVprWjJKcVJuRnVSV0ZsUzNoMlNVTTRTVWNLTUROcmJrODFiVVJuVFdSQ1luQXpXVkJZYVZGV1MyazVSVkpaU0dGbVpEUnllVzl2U25CblQxUkNjQ3N4VmsxMk56WlBRMEptUVhkbloxaHpUVUUwUndwQk1WVmtSSGRGUWk5M1VVVkJkMGxJWjBSQlZFSm5UbFpJVTFWRlJFUkJTMEpuWjNKQ1owVkdRbEZqUkVGNlFXUkNaMDVXU0ZFMFJVWm5VVlZpVjA5UENtMTBkRFV6TW1WTU9WVnRRbFJCVDNkNU4xZFRabk5OZDBoM1dVUldVakJxUWtKbmQwWnZRVlV6T1ZCd2VqRlphMFZhWWpWeFRtcHdTMFpYYVhocE5Ga0tXa1E0ZDJGM1dVUldVakJTUVZGSUwwSkhSWGRZTkZwa1lVaFNNR05JVFRaTWVUbHVZVmhTYjJSWFNYVlpNamwwVERKb01Wb3laSEJpYldSdFdWZE9iQXBNTW5Sc1kyMDFiR0pJVFhSWk1qbDBZbGhXZFdGWVVqVk1lVFZ1WVZoU2IyUlhTWFprTWpsNVlUSmFjMkl6WkhwTU1rb3hZVmQ0YTB4dWJHaGlWM2hCQ21OdFZtMWplVGx2V2xkR2EyTjVPWFJaVjJ4MVRVUnJSME5wYzBkQlVWRkNaemM0ZDBGUlJVVkxNbWd3WkVoQ2VrOXBPSFprUnpseVdsYzBkVmxYVGpBS1lWYzVkV041Tlc1aFdGSnZaRmRLTVdNeVZubFpNamwxWkVkV2RXUkROV3BpTWpCM1NIZFpTMHQzV1VKQ1FVZEVkbnBCUWtGblVWSmtNamw1WVRKYWN3cGlNMlJtV2tkc2VtTkhSakJaTW1kM1RtZFpTMHQzV1VKQ1FVZEVkbnBCUWtGM1VXOVBSMFV6VFhwQ2EwOVVXbXBOZW1NeFRtcENhbGt5V1hoYVJFNXNDazFFYkcxT01rcHBXVEpPYTFsNlp6Uk9hbWQ0VDBkWmVrMTZRVlJDWjI5eVFtZEZSVUZaVHk5TlFVVkZRa0ZXUTJSWGJITmFSRUZ5UW1kdmNrSm5SVVVLUVZsUEwwMUJSVVpDUWpGdlpGZGtibUZYTlc1YWJVWnFXbE01Y2xwWVNuVmFWM2g2VEZkT2RtSlhNVEZpYld3d1pWUkJaRUpuYjNKQ1owVkZRVmxQTHdwTlFVVkhRa0U1ZVZwWFducE1NbWhzV1ZkU2Vrd3lNV2hoVnpSM1QzZFpTMHQzV1VKQ1FVZEVkbnBCUWtOQlVYUkVRM1J2WkVoU2QyTjZiM1pNTTFKMkNtRXlWblZNYlVacVpFZHNkbUp1VFhWYU1td3dZVWhXYVdSWVRteGpiVTUyWW01U2JHSnVVWFZaTWpsMFRVY3dSME5wYzBkQlVWRkNaemM0ZDBGUmEwVUtXSGQ0WkdGSVVqQmpTRTAyVEhrNWJtRllVbTlrVjBsMVdUSTVkRXd5YURGYU1tUndZbTFrYlZsWFRteE1NblJzWTIwMWJHSklUWFJaTWpsMFlsaFdkUXBoV0ZJMVRIazFibUZZVW05a1YwbDJaREk1ZVdFeVduTmlNMlI2VERKS01XRlhlR3RNYm14b1lsZDRRV050Vm0xamVUbHZXbGRHYTJONU9YUlpWMngxQ2sxRVowZERhWE5IUVZGUlFtYzNPSGRCVVc5RlMyZDNiMDlIUlROTmVrSnJUMVJhYWsxNll6Rk9ha0pxV1RKWmVGcEVUbXhOUkd4dFRqSkthVmt5VG1zS1dYcG5ORTVxWjNoUFIxbDZUWHBCWWtKbmIzSkNaMFZGUVZsUEwwMUJSVXhDUVRCTlF6Tk9iR0pIV1hSaFJ6bDZaRWRXYTAxRlFVZERhWE5IUVZGUlFncG5OemgzUVZGM1JVMW5kM2RoU0ZJd1kwaE5Oa3g1T1c1aFdGSnZaRmRKZFZreU9YUk1NbWd4V2pKa2NHSnRaRzFaVjA1c1RESjBiR050Tld4aVNFMTBDbGt5T1hSaVdGWjFZVmhTTlUxRVowZERhWE5IUVZGUlFtYzNPSGRCVVRCRlMyZDNiMDlIUlROTmVrSnJUMVJhYWsxNll6Rk9ha0pxV1RKWmVGcEVUbXdLVFVSc2JVNHlTbWxaTWs1cldYcG5ORTVxWjNoUFIxbDZUWHBCWmtKbmIzSkNaMFZGUVZsUEwwMUJSVTlDUWtWTlJETktiRnB1VFhaaFIxWm9Xa2hOZGdwaVYwWndZbXBCWVVKbmIzSkNaMFZGUVZsUEwwMUJSVkJDUVhkTlEycEZkMDU2UlRCT2VsVXhUV3ByZDB4bldVdExkMWxDUWtGSFJIWjZRVUpGUVZGbkNrUkNOVzlrU0ZKM1kzcHZka3d5WkhCa1IyZ3hXV2sxYW1JeU1IWmhTRlp1V2pKc2RWb3lXbWhaTWxWM1IwRlpTMHQzV1VKQ1FVZEVkbnBCUWtWUlVVc0tSRUZuZVU1VVkzbE5SR013VFhwQ2RFSm5iM0pDWjBWRlFWbFBMMDFCUlZOQ1JqaE5XRmRvTUdSSVFucFBhVGgyV2pKc01HRklWbWxNYlU1MllsTTVid3BrVjJSdVlWYzFibHB0Um1wYVV6bHlXbGhLZFZwWGVIcE1WMDUyWWxjeE1XSnRiREJsVXpoMVdqSnNNR0ZJVm1sTU0yUjJZMjEwYldKSE9UTmplVGxwQ21SWGJITmFRelUxV1ZjeGMxRklTbXhhYmsxMllVZFdhRnBJVFhaaVYwWndZbXBCTkVKbmIzSkNaMFZGUVZsUEwwMUJSVlJDUTI5TlMwUm9hRTU2VFhjS1drUnJNbGw2VFROT1ZGbDNXVEpPYlUxWFVYcGFWRUUxV21wa2FWbHRUbXBhUjAwMFQwUlpORTFVYUcxTmVrMTNTVkZaUzB0M1dVSkNRVWRFZG5wQlFncEdRVkZVUkVKR00ySXpTbkphYlhoMlpERTVhMkZZVG5kWldGSnFZVVJDYTBKbmIzSkNaMFZGUVZsUEwwMUJSVlpDUmxsTlZrZG9NR1JJUW5wUGFUaDJDbG95YkRCaFNGWnBURzFPZG1KVE9XOWtWMlJ1WVZjMWJscHRSbXBhVXpseVdsaEtkVnBYZUhwTVYwNTJZbGN4TVdKdGJEQmxVemxvV1ROU2NHSXlOWG9LVEROS01XSnVUWFpOYW1jMVRYcEZlVTFFU1RCTmVrbDJXVmhTTUZwWE1YZGtTRTEyVFZSQlYwSm5iM0pDWjBWRlFWbFBMMDFCUlZkQ1FXZE5RbTVDTVFwWmJYaHdXWHBDUjBKbmIzSkNaMFZGUVZsUEwwMUJSVmxDUkdkTlRtNUtiR05IT0RaaFNGWnVXakpzZFZveVdtaFpNbFYyWVRKV2VXSnRWbk5qZVRGcUNtSXlNWFJrVnpWd1pFaHJObU50Vm0xUGJrcHNXbTVOZG1GSFZtaGFTRTEyWWxkR2NHSnFRMEpwWjFsTFMzZFpRa0pCU0ZkbFVVbEZRV2RTT0VKSWIwRUtaVUZDTWtGT01EbE5SM0pIZUhoRmVWbDRhMlZJU214dVRuZExhVk5zTmpRemFubDBMelJsUzJOdlFYWkxaVFpQUVVGQlFtNHdSakpaVWpCQlFVRlJSQXBCUldOM1VsRkpaMDlLVjFORmRVOTVjVXc1YmpGU1pFMWlNMEZTWkhGVVNtbFBUVE5JWVV3d2JVbGtZWGN4YTBwaGMydERTVkZEVlhWQmNEVnVia1Z2Q2s5UU0xTlhha1p3ZEZGSVVuWXpjblZzWkRZdlMySk1iamxFZEVWTFZFUnNRMFJCUzBKblozRm9hMnBQVUZGUlJFRjNUbTVCUkVKclFXcENibE15VDJjS05rWmFTVkZuUjBNM1lXdEhZa0ZDVjNKTk1rRnlhV0ZWZVZWbFFUVnJNR05ZWkVSMFZESkZNREpoYTNKbVlXaEtRMWRZUkdKWFNrUkxUbmREVFVjNWNRcFZZekoyUWxSSVkydERUVEJJZDNSMWRFaHhOVWRET1dOUlJrSndNalpyZWpSVlRucFJlVWxVWkZwUWFscFBNamxxVldOM01GSkJTVnA0TUM5aFVUMDlDaTB0TFMwdFJVNUVJRU5GVWxSSlJrbERRVlJGTFMwdExTMEsifX19fQ=="}],"timestampVerificationData":{"rfc3161Timestamps":[{"signedTimestamp":"MIICyTADAgEAMIICwAYJKoZIhvcNAQcCoIICsTCCAq0CAQMxDTALBglghkgBZQMEAgEwgbgGCyqGSIb3DQEJEAEEoIGoBIGlMIGiAgEBBgkrBgEEAYO/MAIwMTANBglghkgBZQMEAgEFAAQgnq4dEiopF6mgNbrNwa0DhBJGjX5TuFh2pqmS2Zm65WsCFQDY9Mr00Cl7hQSo8+vEox9Gn0lPuhgPMjAyNjA3MDgxMTIxNDVaMAMCAQGgMqQwMC4xFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEVMBMGA1UEAxMMc2lnc3RvcmUtdHNhoAAxggHaMIIB1gIBATBRMDkxFTATBgNVBAoTDHNpZ3N0b3JlLmRldjEgMB4GA1UEAxMXc2lnc3RvcmUtdHNhLXNlbGZzaWduZWQCFDoTVC8MkGHuvMFDL8uKjosqI4sMMAsGCWCGSAFlAwQCAaCB/DAaBgkqhkiG9w0BCQMxDQYLKoZIhvcNAQkQAQQwHAYJKoZIhvcNAQkFMQ8XDTI2MDcwODExMjE0NVowLwYJKoZIhvcNAQkEMSIEIFuGEqrjuITTvWRr2FEom6FyD6aYTrAABJKaFL2oiNFiMIGOBgsqhkiG9w0BCRACLzF/MH0wezB5BCCF+Se8B6tiysO0Q1bBDvyBssaIP9p6uebYcNnROs0FtzBVMD2kOzA5MRUwEwYDVQQKEwxzaWdzdG9yZS5kZXYxIDAeBgNVBAMTF3NpZ3N0b3JlLXRzYS1zZWxmc2lnbmVkAhQ6E1QvDJBh7rzBQy/Lio6LKiOLDDAKBggqhkjOPQQDAgRmMGQCMDujXZhHShQhjKUmklOQiWCvm9JXE3+FttJBVkDNsURbwINjnEdzKVb6QQvpLOzwMAIwaTYSmYGVJKszDC7u5XfdojAD2tryUOwiFYv6w5rjwMWjD8irXOCRwNQkoya1HNaw"}]}},"messageSignature":{"messageDigest":{"algorithm":"SHA2_256","digest":"yjffF95tOtpTIBJBWmtMiUUvr3SuvPM9Z5qNg/W4T2U="},"signature":"MEQCIAldB1fO2kXT5hL0Ad/b8dBh1Al80p3EY7aCabkt7zWIAiB6BvpQRvgPmFgc+g5vX5AK6fJeBRlPiNy4WeGZ87hUdg=="}}