| import os |
| import sys |
|
|
| |
| with open(sys.argv[0], 'r') as f: |
| code = f.read() |
| with open(os.path.join(os.path.dirname(sys.argv[0]), 'triton_kernels.py'), 'r') as f: |
| code += f"\n\n{'-'*40}\n# triton_kernels.py\n{'-'*40}\n\n" |
| code += f.read() |
|
|
| import copy |
| import glob |
| import math |
| import threading |
| import time |
| import uuid |
| from dataclasses import dataclass |
| from itertools import accumulate, pairwise |
| from pathlib import Path |
| import gc |
|
|
| os.environ["PYTORCH_ALLOC_CONF"] = "expandable_segments:True" |
| import torch |
| import triton |
| import numpy as np |
|
|
| torch.empty( |
| 1, device=f"cuda:{os.environ['LOCAL_RANK']}", requires_grad=True |
| ).backward() |
| import torch._dynamo as dynamo |
| import torch.distributed as dist |
| import torch.nn.functional as F |
|
|
| |
| from kernels import get_kernel |
| from torch import Tensor, nn |
|
|
| from triton_kernels import XXT, XTX, ba_plus_cAA, FusedLinearReLUSquareFunction, FusedSoftcappedCrossEntropy, transpose_add, transpose_copy |
| |
| |
| ReLUSqrdMLP = FusedLinearReLUSquareFunction.apply |
|
|
| dynamo.config.recompile_limit = 64 |
|
|
| |
| |
| rank = int(os.environ["RANK"]) |
| world_size = int(os.environ["WORLD_SIZE"]) |
| assert 8 % world_size == 0, "world_size must be a divisor of 8" |
| grad_accum_steps = 8 // world_size |
| grad_scale = 1 / grad_accum_steps |
| assert torch.cuda.is_available() |
| device = torch.device("cuda", int(os.environ["LOCAL_RANK"])) |
| torch.cuda.set_device(device) |
| dist.init_process_group(backend="cuda:nccl,cpu:gloo", device_id=device) |
| dist.barrier() |
| master_process = (rank == 0) |
|
|
| |
| |
| |
|
|
| @torch.library.custom_op("nanogpt::mm_t", mutates_args=()) |
| def mm_t_op(x: Tensor, w: Tensor, x_s: float, w_s: float, grad_s: float) -> tuple[Tensor, Tensor, Tensor]: |
| """Computes y = x @ w with F8 weights stored as (in_features, out_features).""" |
| @torch.compile |
| def impl(x: Tensor, w: Tensor): |
| assert x.is_contiguous() and w.is_contiguous() |
| assert x.shape[1] == w.shape[0] |
|
|
| x_f8 = x.div(x_s).to(torch.float8_e4m3fn) |
| w_f8 = w.div(w_s).to(torch.float8_e4m3fn) |
|
|
| |
| |
| w_f8_col_major = w_f8.T.contiguous().T |
|
|
| out = torch._scaled_mm( |
| x_f8, |
| w_f8_col_major, |
| out_dtype=torch.bfloat16, |
| scale_a=x.new_tensor(x_s, dtype=torch.float32), |
| scale_b=x.new_tensor(w_s, dtype=torch.float32), |
| use_fast_accum=True, |
| ) |
| return out, x_f8, w_f8 |
|
|
| return impl(x, w) |
|
|
| @mm_t_op.register_fake |
| def _(x: Tensor, w: Tensor, *_): |
| assert x.ndim == w.ndim == 2 |
| assert x.shape[1] == w.shape[0] |
| assert x.device == w.device |
| assert x.is_contiguous() and w.is_contiguous() |
| return x @ w, x.to(torch.float8_e4m3fn), w.to(torch.float8_e4m3fn) |
|
|
| @torch.library.custom_op("nanogpt::mm_t_backward", mutates_args=()) |
| def mm_t_backward_op(g: Tensor, x_f8: Tensor, w_f8: Tensor, x_s: float, w_s: float, grad_s: float) -> tuple[Tensor, Tensor]: |
| @torch.compile |
| def impl(grad: Tensor, x_f8: Tensor, w_f8: Tensor): |
| assert grad.is_contiguous() |
|
|
| x_scale = grad.new_tensor(x_s, dtype=torch.float32) |
| w_scale = grad.new_tensor(w_s, dtype=torch.float32) |
| grad_scale = grad.new_tensor(grad_s, dtype=torch.float32) |
| grad_f8 = grad.div(grad_s).to(torch.float8_e5m2) |
|
|
| |
| grad_x = torch._scaled_mm( |
| grad_f8, |
| w_f8.T, |
| out_dtype=torch.bfloat16, |
| scale_a=grad_scale, |
| scale_b=w_scale, |
| use_fast_accum=False, |
| ) |
|
|
| |
| |
| grad_w = torch._scaled_mm( |
| x_f8.T.contiguous(), |
| grad_f8.T.contiguous().T, |
| out_dtype=torch.float32, |
| scale_a=x_scale, |
| scale_b=grad_scale, |
| use_fast_accum=False, |
| ) |
|
|
| return grad_x, grad_w |
|
|
| grad_x, grad_w = impl(g, x_f8, w_f8) |
|
|
| return grad_x, grad_w |
|
|
| @mm_t_backward_op.register_fake |
| def _(g: Tensor, x_f8: Tensor, w_f8: Tensor, *_): |
| return x_f8.to(torch.bfloat16), w_f8.to(torch.float32) |
|
|
| def backward_t(ctx, grad_out: Tensor, *_): |
| x_f8, w_f8 = ctx.saved_tensors |
| x_s, w_s, grad_s = ctx.scales |
| grad_x, grad_w = torch.ops.nanogpt.mm_t_backward( |
| grad_out, x_f8, w_f8, x_s, w_s, grad_s |
| ) |
| return grad_x, grad_w, None, None, None |
|
|
| def setup_context_t(ctx: torch.autograd.function.FunctionCtx, inputs, output): |
| *_, x_s, w_s, grad_s = inputs |
| _, x_f8, w_f8 = output |
| ctx.save_for_backward(x_f8, w_f8) |
| ctx.scales = x_s, w_s, grad_s |
| ctx.set_materialize_grads(False) |
|
|
| mm_t_op.register_autograd(backward_t, setup_context=setup_context_t) |
|
|
| |
| |
|
|
| |
| polar_express_coeffs = [ |
| (8.156554524902461, -22.48329292557795, 15.878769915207462), |
| (4.042929935166739, -2.808917465908714, 0.5000178451051316), |
| (3.8916678022926607, -2.772484153217685, 0.5060648178503393), |
| (3.285753657755655, -2.3681294933425376, 0.46449024233003106), |
| (2.3465413258596377, -1.7097828382687081, 0.42323551169305323) |
| ] |
|
|
| @torch.compile(dynamic=False, fullgraph=True) |
| def polar_express(grad_chunk: torch.Tensor, momentum_buffer: torch.Tensor, momentum_t: torch.Tensor, |
| split_baddbmm: bool = False): |
| """ |
| Fused Nesterov momentum + Polar Express Sign Method. |
| Nesterov momentum is applied in FP32, then the result is cast to BF16 for polar express |
| orthogonalization, avoiding materialization of the FP32 intermediate between graph breaks. |
| |
| Polar Express: https://arxiv.org/pdf/2505.16932 |
| by Noah Amsel, David Persson, Christopher Musco, Robert M. Gower. |
| |
| momentum_t is a 0-D CPU tensor to avoid triggering graph recompilations when the value changes. |
| """ |
| |
| momentum = momentum_t.to(grad_chunk.dtype) |
| momentum_buffer.lerp_(grad_chunk, 1 - momentum) |
| g = grad_chunk.lerp_(momentum_buffer, momentum) |
|
|
| X = g.bfloat16() |
| is_tall = g.size(-2) > g.size(-1) |
|
|
| |
| X = X / (X.norm(dim=(-2, -1), keepdim=True) * (1 + 2e-2) + 1e-6) |
|
|
| X = X.contiguous() |
|
|
| if is_tall: |
| |
| A = torch.empty((*X.shape[:-2], X.size(-1), X.size(-1)), device=X.device, dtype=X.dtype) |
| B = torch.empty_like(A) |
| C = torch.empty_like(X) |
|
|
| |
| if split_baddbmm: |
| XB_matmul = torch.bmm if X.ndim > 2 else torch.mm |
| else: |
| aX_plus_XB = torch.baddbmm if X.ndim > 2 else torch.addmm |
|
|
| |
| for a, b, c in polar_express_coeffs: |
| XTX(X, out=A) |
| ba_plus_cAA(A, alpha=c, beta=b, out=B) |
|
|
| |
| |
| |
| |
| if split_baddbmm: |
| XB_matmul(X, B, out=C) |
| C.add_(X, alpha=a) |
| else: |
| aX_plus_XB(X, X, B, beta=a, out=C) |
|
|
| X, C = C, X |
| else: |
| |
| A = torch.empty((*X.shape[:-1], X.size(-2)), device=X.device, dtype=X.dtype) |
| B = torch.empty_like(A) |
| C = torch.empty_like(X) |
|
|
| |
| if split_baddbmm: |
| BX_matmul = torch.bmm if X.ndim > 2 else torch.mm |
| else: |
| aX_plus_BX = torch.baddbmm if X.ndim > 2 else torch.addmm |
|
|
| |
| for a, b, c in polar_express_coeffs: |
| XXT(X, out=A) |
| ba_plus_cAA(A, alpha=c, beta=b, out=B) |
|
|
| if split_baddbmm: |
| BX_matmul(B, X, out=C) |
| C.add_(X, alpha=a) |
| else: |
| aX_plus_BX(X, B, X, beta=a, out=C) |
|
|
| X, C = C, X |
|
|
| return X |
|
|
| |
| |
| def _sparse_comms_active(): |
| |
| return world_size == 8 and grad_accum_steps == 1 |
|
|
| @torch.no_grad |
| def sparse_comms_start(idxes_np, N, rank, world, send_idxes_buffer): |
| rows_per_rank = N // world |
|
|
| |
| send_idxes = send_idxes_buffer[:idxes_np.shape[0]] |
| send_idxes.copy_(torch.from_numpy(idxes_np)) |
| send_idxes = send_idxes.to(device, non_blocking=True) |
|
|
| |
| insertion_points = np.searchsorted( |
| idxes_np, |
| np.arange(0, rows_per_rank * (world + 1), rows_per_rank, dtype=np.int32), |
| ) |
| send_counts = torch.from_numpy(insertion_points[1:] - insertion_points[:-1]) |
| |
| |
| send_counts[rank] = 0 |
|
|
| |
| send_idxes = torch.cat([send_idxes[: insertion_points[rank]], send_idxes[insertion_points[rank + 1] :]]) |
|
|
| |
| |
| recv_counts = torch.empty_like(send_counts) |
| recv_counts_fut = dist.all_to_all_single(recv_counts, send_counts, async_op=True).get_future() |
| return send_idxes, send_counts, recv_counts, recv_counts_fut |
|
|
| @torch.no_grad |
| def sparse_comms_share_indexes(send_idxes, send_counts, recv_counts): |
| |
| total_recv_count = recv_counts.sum().item() |
| recv_counts = recv_counts.tolist() |
| send_counts = send_counts.tolist() |
|
|
| |
| recv_idxes = torch.empty(total_recv_count, dtype=torch.int32, device=device) |
| idxes_fut = dist.all_to_all_single( |
| recv_idxes, |
| send_idxes, |
| output_split_sizes=recv_counts, |
| input_split_sizes=send_counts, |
| async_op=True, |
| ).get_future() |
|
|
| sparse_state = { |
| "send_idxes": send_idxes, |
| "send_counts": send_counts, |
| "recv_counts": recv_counts, |
| } |
| return recv_idxes, sparse_state, idxes_fut |
|
|
| @torch.compile |
| @torch.no_grad |
| def sparse_comms_share_gradients(grad, idxes, send_counts, recv_counts): |
| |
| send_vals = grad[idxes] |
|
|
| d = grad.shape[1] |
|
|
| send_sizes = [i*d for i in send_counts] |
| recv_sizes = [i*d for i in recv_counts] |
|
|
| recv_vals = torch.empty(sum(recv_sizes), device=send_vals.device, dtype=grad.dtype) |
|
|
| val_fut = dist.all_to_all_single( |
| recv_vals, |
| send_vals.view(-1), |
| input_split_sizes=send_sizes, |
| output_split_sizes=recv_sizes, |
| async_op=True, |
| ).get_future() |
|
|
| return recv_vals, val_fut |
|
|
| @torch.no_grad |
| def sparse_comms_merge_gradients(grad, recv_idx, recv_vals, rank, world): |
| d = grad.shape[1] |
| rows_per_rank = grad.shape[0] // world |
|
|
| grad.index_add_(0, recv_idx, recv_vals.view(-1, d)) |
|
|
| |
| return grad[rows_per_rank * rank : rows_per_rank * (rank + 1)].mul_((1 / world)) |
|
|
|
|
| |
| |
|
|
| @dataclass(slots=True) |
| class ParamConfig: |
| """Per-parameter configuration for NorMuonAndAdam optimizer.""" |
| label: str |
| optim: str |
| comms: str |
| adam_betas: tuple[float, float] | None |
| lr_mul: float |
| wd_mul: float |
| lr: float |
| initial_lr: float |
| weight_decay: float |
| |
| eps: float | None = None |
| |
| reshape: tuple | None = None |
| chunk_size: int | None = None |
| momentum: float | None = None |
| beta2: float | None = None |
| per_matrix_lr_mul: list[float] | None = None |
|
|
|
|
| class NorMuonAndAdam: |
| """ |
| Combined optimizer that handles both NorMuon (for projection matrices) and |
| Adam (for embeddings/scalars/gate weights). |
| |
| Muon - MomentUm Orthogonalized by Newton-schulz |
| |
| https://kellerjordan.github.io/posts/muon/ |
| |
| Muon internally runs standard SGD-momentum, and then performs an orthogonalization post- |
| processing step, in which each 2D parameter's update is replaced with the nearest orthogonal |
| matrix. To efficiently orthogonalize each update, Muon uses a Newton-Schulz iteration (replaced |
| here with Polar Express), which has the advantage that it can be stably run in bfloat16 on the GPU. |
| |
| Muon is applied only to the projection matrices in the attention and MLP layers, and is not recommended |
| for embeddings, scalars, or individual weight vectors (e.g., bias terms or gate weights). |
| |
| Differences from standard Muon: |
| - Newton-Shulz is replaced with Polar Express for the orthogonalization step |
| - NorMuon adds a low-rank variance estimator similar to Adafactor. https://arxiv.org/pdf/2510.05491 |
| - Cautious weight decay, a gated version of decoupled weight decay |
| - Mantissa tracking for precision |
| |
| Adam (for embeddings/scalars/gates): |
| - Standard Adam with bias correction |
| - Cautious weight decay |
| |
| Configuration: |
| Unlike torch.optim.Optimizer, this class uses per-parameter configs from a `param_table` dict |
| and does not include parameter "groups". All parameters require a .label attribute, and a |
| corresponding entry in the param_table to specify their hyperparameters (lr_mul, wd_mul, adam_betas, etc.). |
| |
| Communication and ordering: |
| Gradient communication is explicitly scheduled rather than hook-driven. |
| Reductions are launched in `scatter_order`, while update math and final |
| gathers are executed in `work_order`. These orders are independent and |
| must each contain every parameter label exactly once. |
| |
| Two communication modes are supported per parameter: |
| - 'replicated': Gradients are all-reduced and each rank computes the full update. |
| - 'sharded': Gradients are reduce-scattered, each rank updates its shard, |
| and results are all-gathered. |
| |
| Adam parameters may be freely sharded. NorMuon operates on full matrices; sharding is |
| supported by grouping matrices into parameter banks. NorMuon parameters must have a |
| `.reshape` attribute that reshapes the bank so that the leading dimension is divisible |
| by world_size. |
| |
| # Contributors include @YouJiacheng, @KonstantinWilleke, @alexrgilbert, @adricarda, |
| # @tuttyfrutyee, @vdlad, @ryanyang0, @vagrawal, @varunneal, @chrisjmccormick |
| """ |
| def __init__(self, named_params, param_table: dict, scatter_order: list, work_order: list, |
| adam_defaults: dict, normuon_defaults: dict): |
| self.world_size = dist.get_world_size() if dist.is_initialized() else 1 |
|
|
| |
| self.adam_defaults = adam_defaults |
| self.normuon_defaults = normuon_defaults |
| self.param_table = param_table |
| self.scatter_order = scatter_order |
| self.work_order = work_order |
|
|
| |
| self.param_cfgs: dict[nn.Parameter, ParamConfig] = {} |
| self.param_states: dict[nn.Parameter, dict] = {} |
| self._param_by_label: dict[str, nn.Parameter] = {} |
| for name, param in named_params: |
| label = getattr(param, "label", None) |
| assert label is not None and label in param_table |
| assert label not in self._param_by_label |
| self._param_by_label[label] = param |
| self._build_param_cfg(param, label) |
|
|
| |
| present = self._param_by_label.keys() |
| assert set(scatter_order) == present and set(work_order) == present |
|
|
| |
| if self.world_size == 1: |
| for p_cfg in self.param_cfgs.values(): |
| p_cfg.comms = "none" |
|
|
| |
| self._init_state() |
|
|
| |
| self._step_size_t = torch.tensor(0.0, dtype=torch.float32, device="cpu") |
| self._eff_wd_t = torch.tensor(0.0, dtype=torch.float32, device="cpu") |
| self._eff_lr_t = torch.tensor(0.0, dtype=torch.float32, device="cpu") |
| self._momentum_t = torch.tensor(0.0, dtype=torch.float32, device="cpu") |
|
|
| |
| self._reduce_futures: dict[nn.Parameter, tuple] = {} |
| self._sparse_async_data: dict[nn.Parameter, list] = {} |
|
|
| |
| self.split_embed = False |
| self._lm_head_param = self._param_by_label.get("lm_head") |
| self._embed_param = self._param_by_label.get("embed") |
|
|
| def _build_param_cfg(self, param: nn.Parameter, label: str): |
| """Build config for a single parameter from param_table.""" |
| table_entry = self.param_table[label] |
| optim = table_entry["optim"] |
| comms = table_entry["comms"] |
| if comms == "sharded_sparse" and not _sparse_comms_active(): |
| comms = "sharded" |
| adam_betas = table_entry.get("adam_betas") |
| lr_mul = table_entry.get("lr_mul", 1.0) |
| wd_mul = table_entry.get("wd_mul", 1.0) |
|
|
| if optim == "adam": |
| chunk_size = param.shape[0] // self.world_size if comms.startswith("sharded") else None |
| p_cfg = ParamConfig( |
| label=label, |
| optim=optim, |
| comms=comms, |
| adam_betas=tuple(adam_betas) if adam_betas else None, |
| lr_mul=lr_mul, |
| wd_mul=wd_mul, |
| lr=self.adam_defaults["lr"], |
| initial_lr=self.adam_defaults["lr"], |
| weight_decay=self.adam_defaults["weight_decay"], |
| eps=self.adam_defaults["eps"], |
| chunk_size=chunk_size, |
| ) |
| elif optim == "normuon": |
| reshape = getattr(param, "reshape", None) |
| if reshape is None: |
| raise ValueError(f"NorMuon param {label} must have .reshape attribute") |
| if reshape[0] % self.world_size != 0: |
| raise ValueError(f"reshape[0]={reshape[0]} must be divisible by world_size") |
|
|
| chunk_size = reshape[0] // self.world_size |
| chunk_shape = (chunk_size, *reshape[1:]) |
| |
| shape_mult = max(1.0, chunk_shape[-2] / chunk_shape[-1]) ** 0.5 if len(chunk_shape) >= 2 else 1.0 |
| lr_mul = shape_mult * lr_mul |
|
|
| |
| per_matrix_lr_mul = None |
| if label == "mlp_bank": |
| rank = dist.get_rank() if dist.is_initialized() else 0 |
| start_idx = rank * chunk_size |
| per_matrix_lr_mul = [] |
| for i in range(chunk_size): |
| global_idx = start_idx + i |
| is_c_proj = (global_idx % 2 == 1) |
| per_matrix_lr_mul.append(2.0 if is_c_proj else 1.0) |
|
|
| p_cfg = ParamConfig( |
| label=label, |
| optim=optim, |
| comms=comms, |
| adam_betas=tuple(adam_betas) if adam_betas else None, |
| lr_mul=lr_mul, |
| wd_mul=wd_mul, |
| lr=self.normuon_defaults["lr"], |
| initial_lr=self.normuon_defaults["lr"], |
| weight_decay=self.normuon_defaults["weight_decay"], |
| reshape=reshape, |
| chunk_size=chunk_size, |
| momentum=self.normuon_defaults["momentum"], |
| beta2=self.normuon_defaults["beta2"], |
| per_matrix_lr_mul=per_matrix_lr_mul, |
| ) |
| else: |
| raise ValueError(f"Unknown optim type: {optim}") |
|
|
| self.param_cfgs[param] = p_cfg |
|
|
| def _init_state(self): |
| """Initialize optimizer state for all parameters.""" |
| for param, p_cfg in self.param_cfgs.items(): |
| if p_cfg.optim == "adam": |
| |
| if p_cfg.comms.startswith("sharded"): |
| chunk = param[:p_cfg.chunk_size] |
| else: |
| chunk = param |
| exp_avg = torch.zeros_like(chunk, dtype=torch.float32, device=param.device) |
| self.param_states[param] = dict(step=0, exp_avg=exp_avg, exp_avg_sq=torch.zeros_like(exp_avg)) |
|
|
| elif p_cfg.optim == "normuon": |
| chunk_shape = (p_cfg.chunk_size, *p_cfg.reshape[1:]) |
|
|
| |
| momentum_buffer = torch.zeros( |
| chunk_shape, dtype=torch.float32, device=param.device |
| ) |
|
|
| |
| if chunk_shape[-2] >= chunk_shape[-1]: |
| second_mom_shape = (*chunk_shape[:-1], 1) |
| else: |
| second_mom_shape = (*chunk_shape[:-2], 1, chunk_shape[-1]) |
| second_momentum_buffer = torch.zeros( |
| second_mom_shape, dtype=torch.float32, device=param.device |
| ) |
|
|
| |
| mantissa = torch.zeros( |
| chunk_shape, dtype=torch.uint16, device=param.device |
| ) |
|
|
| self.param_states[param] = dict( |
| momentum_buffer=momentum_buffer, |
| second_momentum_buffer=second_momentum_buffer, |
| mantissa=mantissa, |
| ) |
|
|
| |
| |
|
|
| def _launch_reduce(self, param: nn.Parameter, grad: Tensor): |
| """Launch async reduce for a parameter based on its comms policy.""" |
| p_cfg = self.param_cfgs[param] |
|
|
| if p_cfg.comms == "none": |
| if p_cfg.optim == "normuon": |
| |
| grad = grad.view(p_cfg.reshape) |
| self._reduce_futures[param] = (None, grad) |
| elif p_cfg.comms == "replicated": |
| future = dist.all_reduce(grad, op=dist.ReduceOp.AVG, async_op=True).get_future() |
| self._reduce_futures[param] = (future, grad) |
| elif p_cfg.comms == "sharded": |
| if p_cfg.optim == "normuon": |
| |
| grad_reshaped = grad.view(p_cfg.reshape) |
| grad_chunk = torch.empty( |
| (p_cfg.chunk_size, *grad_reshaped.shape[1:]), |
| dtype=grad.dtype, |
| device=grad.device |
| ) |
| future = dist.reduce_scatter_tensor( |
| grad_chunk, grad_reshaped.contiguous(), op=dist.ReduceOp.AVG, async_op=True |
| ).get_future() |
| self._reduce_futures[param] = (future, grad_chunk) |
| else: |
| |
| grad_chunk = torch.empty_like(grad[:p_cfg.chunk_size]) |
| future = dist.reduce_scatter_tensor( |
| grad_chunk, grad, op=dist.ReduceOp.AVG, async_op=True |
| ).get_future() |
| self._reduce_futures[param] = (future, grad_chunk) |
| elif p_cfg.comms == "sharded_sparse": |
| sparse_state = self._sparse_async_data[param] |
| send_idxes = sparse_state["send_idxes"] |
| send_counts = sparse_state["send_counts"] |
| recv_counts = sparse_state["recv_counts"] |
| recv_vals, val_fut = sparse_comms_share_gradients( |
| grad, send_idxes, send_counts, recv_counts |
| ) |
| self._reduce_futures[param].extend((val_fut, recv_vals)) |
|
|
| def _launch_gather(self, param: nn.Parameter, p_slice: Tensor) -> "torch.futures.Future": |
| """Launch async all_gather for a sharded parameter.""" |
| p_cfg = self.param_cfgs[param] |
| if p_cfg.optim == "normuon": |
| full_param = param.data.view(p_cfg.reshape) |
| assert full_param.is_contiguous() |
| return dist.all_gather_into_tensor( |
| full_param, p_slice.contiguous(), async_op=True |
| ).get_future() |
| else: |
| return dist.all_gather_into_tensor( |
| param, p_slice.contiguous(), async_op=True |
| ).get_future() |
|
|
| |
| |
|
|
| def reset(self): |
| """Reset NorMuon momentum buffers and split_embed state (called on training reset).""" |
| self.split_embed = False |
| for param, p_cfg in self.param_cfgs.items(): |
| if p_cfg.optim == "normuon": |
| p_state = self.param_states[param] |
| p_state["momentum_buffer"].zero_() |
| p_state["mantissa"].zero_() |
| p_state["second_momentum_buffer"].zero_() |
|
|
| def copy_lm_state_to_embed(self): |
| """ |
| Copy the optimizer state from the lm_head to the embed at the untie point. |
| This requires an all-gather + reshard because of different sharding: |
| - lm_head (768, 50304) is sharded to (96, 50304) per rank (along model_dim) |
| - embed (50304, 768) is sharded to (6288, 768) per rank (along vocab_size) |
| |
| We all-gather the lm_head momentum, transpose it, then each rank takes their |
| embed shard to get the correct momentum state. |
| """ |
| lm_head = self._lm_head_param |
| embed = self._embed_param |
| lm_state = self.param_states[lm_head] |
| embed_state = self.param_states[embed] |
| lm_cfg = self.param_cfgs[lm_head] |
| embed_cfg = self.param_cfgs[embed] |
|
|
| embed_state['step'] = lm_state['step'] |
|
|
| |
| if self.world_size > 1: |
| rank = dist.get_rank() |
| lm_chunk_size = lm_cfg.chunk_size |
| embed_chunk_size = embed_cfg.chunk_size |
|
|
| |
| for key in ["exp_avg", "exp_avg_sq"]: |
| lm_chunk = lm_state[key] |
| full_lm = torch.empty(lm_head.shape[0], lm_head.shape[1], dtype=lm_chunk.dtype, device=lm_chunk.device) |
| dist.all_gather_into_tensor(full_lm, lm_chunk.contiguous()) |
| embed_state[key].copy_(full_lm.T[rank * embed_chunk_size:(rank + 1) * embed_chunk_size]) |
| else: |
| |
| for key in ["exp_avg", "exp_avg_sq"]: |
| embed_state[key].copy_(lm_state[key].T) |
|
|
| |
| self.split_embed = True |
|
|
| def state_dict(self): |
| """Return the optimizer state as a dict.""" |
| return { |
| "param_states": {id(p): s for p, s in self.param_states.items()}, |
| "param_cfgs": {id(p): s for p, s in self.param_cfgs.items()}, |
| } |
|
|
| def load_state_dict(self, state_dict): |
| """Load optimizer state from a dict.""" |
| |
| id_to_param = {id(p): p for p in self.param_cfgs} |
|
|
| |
| for param_id, saved_p_state in state_dict["param_states"].items(): |
| if param_id in id_to_param: |
| param = id_to_param[param_id] |
| p_state = self.param_states[param] |
| for k, v in saved_p_state.items(): |
| if isinstance(v, torch.Tensor) and k in p_state: |
| target_dtype = p_state[k].dtype |
| p_state[k] = v.to(dtype=target_dtype, device=p_state[k].device) |
| else: |
| p_state[k] = v |
|
|
| |
| |
|
|
| @torch.no_grad() |
| def step(self, do_adam: bool = True): |
| """ |
| Combined optimizer step with explicit ordering. |
| |
| Args: |
| do_adam: If True, update Adam params. NorMuon params always updated. |
| |
| Flow: |
| 1. Scatter phase: Launch reduces in scatter_order |
| 2. Work phase: Process updates in work_order |
| - Wait for reduce, compute update, launch gather |
| 3. Finalize phase: Wait for gathers |
| |
| While the embeddings are tied: |
| - Comms and update math are only done on lm_head. |
| - We add embed.grad.T into lm_head.grad before comms. |
| - After lm_head gather, we copy lm_head.data.T --> embed.data |
| """ |
| rank = dist.get_rank() if dist.is_initialized() else 0 |
| lm_param, embed_param = self._lm_head_param, self._embed_param |
|
|
| |
| for label in self.scatter_order: |
| param = self._param_by_label[label] |
| p_cfg = self.param_cfgs[param] |
|
|
| if p_cfg.optim == "adam" and not do_adam: |
| continue |
| if param.grad is None: |
| continue |
|
|
| |
| if label == "lm_head" and do_adam and not self.split_embed: |
| if embed_param is not None and embed_param.grad is not None: |
| transpose_add(embed_param.grad, param.grad) |
|
|
| |
| if label == "embed" and not self.split_embed: |
| continue |
|
|
| self._launch_reduce(param, param.grad) |
|
|
| |
| gather_futures = [] |
| lm_head_gather_future = None |
|
|
| for label in self.work_order: |
| param = self._param_by_label[label] |
| if param not in self._reduce_futures: |
| continue |
|
|
| p_cfg = self.param_cfgs[param] |
| if p_cfg.optim == "adam" and not do_adam: |
| continue |
| |
| if p_cfg.comms != "sharded_sparse": |
| future, grad_chunk = self._reduce_futures[param] |
| if future is not None: |
| future.wait() |
| else: |
| idxes_fut, recv_idxes, recv_fut, recv_vals = self._reduce_futures[param] |
| idxes_fut.wait() |
| recv_fut.wait() |
|
|
| grad_chunk = sparse_comms_merge_gradients(param.grad, recv_idxes, recv_vals, rank, world_size) |
|
|
| |
| if p_cfg.optim == "adam": |
| p_slice = self._adam_update(param, grad_chunk, p_cfg, rank) |
| else: |
| p_slice = self._normuon_update(param, grad_chunk, p_cfg, rank) |
| |
| if p_cfg.comms.startswith("sharded") and self.world_size > 1: |
| gather_fut = self._launch_gather(param, p_slice) |
| if label == "lm_head": |
| lm_head_gather_future = gather_fut |
| else: |
| gather_futures.append(gather_fut) |
|
|
| |
| |
| if lm_head_gather_future is not None: |
| lm_head_gather_future.wait() |
|
|
| |
| if do_adam and not self.split_embed and embed_param is not None and lm_param is not None: |
| transpose_copy(lm_param.data, embed_param.data) |
|
|
| |
| for fut in gather_futures: |
| fut.wait() |
|
|
| self._reduce_futures.clear() |
| self._sparse_async_data.clear() |
|
|
| |
| for param, p_cfg in self.param_cfgs.items(): |
| if p_cfg.optim == "adam" and not do_adam: |
| continue |
| param.grad = None |
|
|
| |
| |
|
|
| def _adam_update(self, param: nn.Parameter, grad_chunk: Tensor, p_cfg: ParamConfig, rank: int) -> Tensor: |
| """Apply Adam update to a parameter. Returns the updated p_slice.""" |
| beta1, beta2 = p_cfg.adam_betas |
| lr = p_cfg.lr * p_cfg.lr_mul |
|
|
| |
| if p_cfg.comms.startswith("sharded"): |
| p_slice = param[rank * p_cfg.chunk_size:(rank + 1) * p_cfg.chunk_size] |
| else: |
| p_slice = param |
|
|
| p_state = self.param_states[param] |
| p_state["step"] += 1 |
| t = p_state["step"] |
|
|
| bias1, bias2 = 1 - beta1 ** t, 1 - beta2 ** t |
| self._step_size_t.fill_(lr * (bias2 ** 0.5 / bias1)) |
| self._eff_wd_t.fill_(lr * lr * p_cfg.weight_decay * p_cfg.wd_mul) |
|
|
| NorMuonAndAdam._adam_update_step( |
| p_slice, grad_chunk, p_state["exp_avg"], p_state["exp_avg_sq"], |
| beta1, beta2, p_cfg.eps, self._step_size_t, self._eff_wd_t |
| ) |
|
|
| return p_slice |
|
|
| @staticmethod |
| @torch.compile(dynamic=False, fullgraph=True) |
| def _adam_update_step(p_slice, g_slice, exp_avg, exp_avg_sq, beta1, beta2, eps, step_size_t, eff_wd_t): |
| """Compiled Adam update step.""" |
| exp_avg.mul_(beta1).add_(g_slice, alpha=1 - beta1) |
| exp_avg_sq.mul_(beta2).addcmul_(g_slice, g_slice, value=1 - beta2) |
| update = exp_avg.div(exp_avg_sq.sqrt().add_(eps)).mul_(step_size_t) |
| |
| mask = (update * p_slice) > 0 |
| update.addcmul_(p_slice, mask, value=eff_wd_t) |
| p_slice.add_(other=update, alpha=-1.0) |
|
|
| |
| |
|
|
| def _normuon_update(self, param: nn.Parameter, grad_chunk: Tensor, p_cfg: ParamConfig, rank: int) -> Tensor: |
| """Apply NorMuon update to a parameter. Returns the updated p_slice.""" |
| chunk_shape = grad_chunk.shape |
|
|
| p_state = self.param_states[param] |
| grad_chunk = grad_chunk.float() |
|
|
| self._momentum_t.fill_(p_cfg.momentum) |
| self._eff_lr_t.fill_(p_cfg.lr_mul * p_cfg.lr) |
| self._eff_wd_t.fill_(p_cfg.wd_mul * p_cfg.weight_decay * p_cfg.lr) |
|
|
| |
| is_large_matrix = chunk_shape[-2] > 1024 |
| v_chunk = polar_express( |
| grad_chunk, p_state["momentum_buffer"], self._momentum_t, |
| split_baddbmm=is_large_matrix, |
| ) |
|
|
| |
| red_dim = -1 if chunk_shape[-2] >= chunk_shape[-1] else -2 |
| v_chunk = NorMuonAndAdam._apply_normuon_variance_reduction( |
| v_chunk, p_state["second_momentum_buffer"], p_cfg.beta2, red_dim |
| ) |
|
|
| |
| param_view = param.data.view(p_cfg.reshape) |
| p_slice = param_view[rank * p_cfg.chunk_size:(rank + 1) * p_cfg.chunk_size] |
|
|
| |
| if p_cfg.per_matrix_lr_mul is not None: |
| self._eff_wd_t.fill_(p_cfg.wd_mul * p_cfg.weight_decay * p_cfg.lr) |
| for mat_idx in range(p_cfg.chunk_size): |
| self._eff_lr_t.fill_(p_cfg.lr_mul * p_cfg.per_matrix_lr_mul[mat_idx] * p_cfg.lr) |
| NorMuonAndAdam._cautious_wd_and_update_inplace( |
| p_slice[mat_idx].view(torch.uint16), p_state["mantissa"][mat_idx], v_chunk[mat_idx], |
| self._eff_wd_t, self._eff_lr_t |
| ) |
| else: |
| NorMuonAndAdam._cautious_wd_and_update_inplace( |
| p_slice.view(torch.uint16), p_state["mantissa"], v_chunk, |
| self._eff_wd_t, self._eff_lr_t |
| ) |
|
|
| return p_slice |
|
|
| @staticmethod |
| @torch.compile(dynamic=False, fullgraph=True) |
| def _cautious_wd_and_update_inplace(p, mantissa, grad, wd_tensor, lr_tensor): |
| """ |
| Cautious weight decay + parameter update. wd_tensor and lr_tensor are 0-D CPU tensors. |
| Mantissa is tracked to enable higher precision updates on bfloat16 parameters. |
| bfloat16 format: 1 sign bit + 8 exponent bits + 7 mantissa bits = 16 bits total |
| float32 format: 1 sign bit + 8 exponent bits + 23 mantissa bits = 32 bits total |
| """ |
| assert p.dtype == mantissa.dtype == torch.uint16 |
| grad = grad.float() |
| wd_factor = wd_tensor.to(torch.float32) |
| lr_factor = lr_tensor.to(torch.float32) |
| p_precise_raw = (p.to(torch.uint32) << 16) | mantissa.to(torch.uint32) |
| p_precise = p_precise_raw.view(torch.float32) |
| mask = (grad * p_precise) >= 0 |
| p_precise.copy_(p_precise - (p_precise * mask * wd_factor * lr_factor) - (grad * lr_factor)) |
| p.copy_((p_precise_raw >> 16).to(torch.uint16)) |
| mantissa.copy_(p_precise_raw.to(torch.uint16)) |
|
|
| @staticmethod |
| @torch.compile(dynamic=False, fullgraph=True) |
| def _apply_normuon_variance_reduction(v_chunk, second_momentum_buffer, beta2, red_dim): |
| """NorMuon variance reduction. Algebraically fuses the normalization steps to minimize memory ops.""" |
| v_mean = v_chunk.float().square().mean(dim=red_dim, keepdim=True) |
| red_dim_size = v_chunk.size(red_dim) |
| v_norm_sq = v_mean.sum(dim=(-2, -1), keepdim=True).mul_(red_dim_size) |
| v_norm = v_norm_sq.sqrt_() |
| second_momentum_buffer.lerp_(v_mean.to(dtype=second_momentum_buffer.dtype), 1 - beta2) |
| step_size = second_momentum_buffer.clamp_min(1e-10).rsqrt_() |
| scaled_sq_sum = (v_mean * red_dim_size) * step_size.float().square() |
| v_norm_new = scaled_sq_sum.sum(dim=(-2, -1), keepdim=True).sqrt_() |
| final_scale = step_size * (v_norm / v_norm_new.clamp_min_(1e-10)) |
| return v_chunk.mul_(final_scale.type_as(v_chunk)) |
|
|
| |
| |
|
|
| def norm(x: Tensor): |
| return F.rms_norm(x, (x.size(-1),)) |
|
|
|
|
| class CastedLinearT(nn.Module): |
| """ |
| Linear layer with transposed weight storage (in_features, out_features) which |
| addresses the slow kernel that was used for gradient accumulation. @chrisjmccormick |
| """ |
| def __init__(self, in_features: int, out_features: int, use_fp8=False, x_s=1.0, w_s=1.0, grad_s=1.0): |
| super().__init__() |
| self.in_features = in_features |
| self.out_features = out_features |
| self.use_fp8 = use_fp8 |
| self.x_s = x_s |
| self.w_s = w_s |
| self.grad_s = grad_s |
|
|
| self.weight = nn.Parameter(torch.empty(in_features, out_features, dtype=torch.bfloat16)) |
| self.reset_parameters() |
|
|
| def reset_parameters(self) -> None: |
| with torch.no_grad(): |
| nn.init.zeros_(self.weight) |
|
|
| def forward(self, x: Tensor): |
| if self.use_fp8 and self.training: |
| _x = x.flatten(0, -2) |
| out = torch.ops.nanogpt.mm_t(_x, self.weight, x_s=self.x_s, w_s=self.w_s, grad_s=self.grad_s)[0] |
| return out.reshape(*x.shape[:-1], -1) |
| else: |
| return x @ self.weight.type_as(x) |
|
|
| |
| |
|
|
| class Yarn(nn.Module): |
| def __init__(self, head_dim, max_seq_len, paired=False): |
| super().__init__() |
| self.head_dim = head_dim |
| self.max_seq_len = max_seq_len |
| self.paired = paired |
| self.reset() |
|
|
| def rotary(self, x_BTHD): |
| assert self.factor1.size(0) >= x_BTHD.size(-3) |
| factor1, factor2 = ( |
| self.factor1[None, : x_BTHD.size(-3), None, :], |
| self.factor2[None, : x_BTHD.size(-3), None, :], |
| ) |
| x_flip = x_BTHD.view(*x_BTHD.shape[:-1], x_BTHD.shape[-1] // 2, 2).flip(-1).view(x_BTHD.shape) |
| return factor1 * x_BTHD + factor2 * x_flip |
|
|
| def reset(self): |
| angular_freq = (1 / 1024) ** torch.linspace(0, 1, steps=self.head_dim//4, dtype=torch.float32, device=device) |
| angular_freq = angular_freq.repeat_interleave(2) |
| |
| angular_freq = torch.cat([angular_freq, angular_freq.new_zeros(self.head_dim//2)]) |
| t = torch.arange(2*self.max_seq_len, dtype=torch.float32, device=device) |
| if not self.paired: |
| theta = torch.outer(t, angular_freq) |
| self.factor1 = nn.Buffer( |
| theta.cos().to(torch.bfloat16), persistent=False |
| ) |
| self.factor2 = nn.Buffer( |
| theta.sin().to(torch.bfloat16), persistent=False |
| ) |
| else: |
| t_even = 2 * t |
| t_odd = t_even + 1 |
| theta1 = torch.outer(t_even, angular_freq) |
| theta2 = torch.outer(t_odd, angular_freq) |
| self.factor1 = nn.Buffer( |
| torch.cat((theta1.cos(), theta2.cos()), dim=-1).to(torch.bfloat16), |
| persistent=False |
| ) |
| self.factor2 = nn.Buffer( |
| torch.cat((theta1.sin(), theta2.sin()), dim=-1).to(torch.bfloat16), |
| persistent=False |
| ) |
| self.factor2[..., 1::2] *= -1 |
| self.angular_freq = angular_freq |
| |
| self.attn_scale = 0.1 |
|
|
| def apply(self, old_window: int, new_window: int, alpha: int=1, beta: int=32): |
| rotations = old_window * self.angular_freq / (2 * torch.pi) |
| scaling_factor = old_window / new_window |
| interpolation_weight = torch.clamp((rotations - alpha) / (beta - alpha), 0, 1) |
| self.angular_freq *= scaling_factor + interpolation_weight * (1 - scaling_factor) |
| t = torch.arange(2*self.max_seq_len, dtype=torch.float32, device=self.angular_freq.device) |
| if not self.paired: |
| theta = torch.outer(t, self.angular_freq) |
| self.factor1.copy_(theta.cos()) |
| self.factor2.copy_(theta.sin()) |
| else: |
| t_even = 2 * t |
| t_odd = t_even + 1 |
| theta1 = torch.outer(t_even, self.angular_freq) |
| theta2 = torch.outer(t_odd, self.angular_freq) |
| self.factor1.copy_(torch.cat((theta1.cos(), theta2.cos()), dim=-1)) |
| self.factor2.copy_(torch.cat((theta1.sin(), theta2.sin()), dim=-1)) |
| self.factor2[..., 1::2] *= -1 |
| self.attn_scale *= 0.2 * math.log(new_window / old_window) + 1 |
|
|
| @dataclass(slots=True) |
| class AttnArgs: |
| sa_lambdas: torch.Tensor |
| seqlens: torch.Tensor |
| bm_size: int |
| yarn: Yarn |
| key_offset: bool |
| attn_gate_w: torch.Tensor |
| aux_v: torch.Tensor | None |
| xsa_alpha: torch.Tensor | None |
| train_max_seq_len: torch.Tensor |
|
|
| flash_attn_interface = get_kernel('kernels-community/flash-attn3', version=1).flash_attn_interface |
|
|
| class CausalSelfAttention(nn.Module): |
| def __init__(self, dim: int, head_dim: int, num_heads: int, paired: bool = False): |
| super().__init__() |
| self.num_heads = num_heads |
| self.head_dim = head_dim |
| self.dim = dim |
| self.hdim = num_heads * head_dim |
| self.paired = paired |
| assert self.hdim == self.dim, "num_heads * head_dim must equal model_dim" |
| |
|
|
| def forward(self, x: Tensor, attn_args: AttnArgs, qkvo_w: Tensor): |
| B, T = x.size(0), x.size(1) |
| assert B == 1, "varlen sequences requires B == 1" |
| assert T % 16 == 0 |
| |
| aux_v, attn_gate_w = attn_args.aux_v, attn_args.attn_gate_w |
| sa_lambdas, key_offset = attn_args.sa_lambdas, attn_args.key_offset |
| seqlens, bm_size = attn_args.seqlens, attn_args.bm_size |
| train_max_seq_len, yarn = attn_args.train_max_seq_len, attn_args.yarn |
|
|
| q, k, v = F.linear(x, sa_lambdas[0] * qkvo_w[:self.dim * 3].type_as(x)).view(B, T, 3 * self.num_heads, self.head_dim).chunk(3, dim=-2) |
| max_len = train_max_seq_len if self.training else (args.val_batch_size // (grad_accum_steps * world_size)) |
|
|
| q, k = norm(q), norm(k) |
|
|
| if not self.paired: |
| q, k = yarn.rotary(q), yarn.rotary(k) |
|
|
| if key_offset: |
| |
| k[:, 1:, :, self.head_dim // 2:] = k[:, :-1, :, self.head_dim // 2:] |
|
|
| if aux_v is not None: |
| v = v + aux_v.view_as(v) |
|
|
| else: |
| |
| |
| |
| |
| q = q.view(B, T, self.num_heads // 2, self.head_dim * 2) |
| k = k.view(B, T, self.num_heads // 2, self.head_dim * 2) |
| v = v.reshape(B, T * 2, self.num_heads // 2, self.head_dim) |
|
|
| q, k = yarn.rotary(q), yarn.rotary(k) |
|
|
| q = q.view(B, T * 2, self.num_heads // 2, self.head_dim) |
| k = k.view(B, T * 2, self.num_heads // 2, self.head_dim) |
|
|
| if aux_v is not None: |
| v = v + aux_v.view_as(v) |
|
|
| seqlens = 2 * seqlens |
| max_len = 2 * max_len |
|
|
| |
| y = flash_attn_interface.flash_attn_varlen_func(q[0], k[0], v[0], cu_seqlens_q=seqlens, cu_seqlens_k=seqlens, |
| max_seqlen_q=max_len, max_seqlen_k=max_len, |
| causal=True, softmax_scale=yarn.attn_scale, window_size=(bm_size, 0)) |
| y = y.view(B, T, self.num_heads, self.head_dim) |
| |
| |
| if attn_args.xsa_alpha is not None and not self.paired: |
| vn = F.normalize(v, dim=-1, eps=1e-4) |
| proj = (y * vn).sum(-1, keepdim=True) |
| alpha = torch.tanh(attn_args.xsa_alpha).type_as(y).view(1, 1, self.num_heads, 1) |
| y = y - alpha * proj * vn |
| y = y * torch.sigmoid(F.linear(x[..., :12], attn_gate_w)).view(B, T, self.num_heads, 1) |
| y = y.contiguous().view(B, T, self.num_heads * self.head_dim) |
| y = F.linear(y, sa_lambdas[1] * qkvo_w[self.dim * 3:].type_as(y)) |
| return y |
|
|
|
|
| |
| |
|
|
| def next_multiple_of_n(v: float | int, *, n: int): |
| return math.ceil(v / n) * n |
|
|
| @dataclass(slots=True) |
| class ForwardScheduleConfig: |
| mtp_weights: torch.Tensor |
| ws_short: int |
| ws_long: int |
| train_max_seq_len: int |
|
|
| class GPT(nn.Module): |
| def __init__(self, vocab_size: int, num_layers: int, num_heads: int, head_dim: int, model_dim: int, max_seq_len: int): |
| super().__init__() |
| self.num_layers = num_layers |
| self.num_heads = num_heads |
| self.head_dim = head_dim |
| |
| |
| self.vocab_size = next_multiple_of_n(vocab_size, n=128) |
|
|
| |
| use_fp8 = not os.environ.get("DISABLE_FP8", False) |
| self.lm_head = CastedLinearT(model_dim, self.vocab_size, use_fp8=use_fp8, x_s=100/448, w_s=1.6/448, grad_s=grad_scale * 0.75/448) |
| nn.init.normal_(self.lm_head.weight, mean=0, std=0.005) |
|
|
| self.embed = nn.Embedding(self.vocab_size, model_dim) |
| with torch.no_grad(): |
| |
| self.embed.weight.copy_(self.lm_head.weight.T) |
|
|
| self.init_attn(model_dim, head_dim, num_heads, num_layers, max_seq_len) |
| self.init_mlp(model_dim) |
| self.init_misc(model_dim, num_layers) |
| self.init_mudd(num_layers, model_dim) |
|
|
| |
| for name, param in self.named_parameters(): |
| param.label = name.replace('.weight', '') |
|
|
| def init_attn(self, model_dim, head_dim, num_heads, num_layers, max_seq_len): |
| |
| self.cache_layers = [3, 7] |
|
|
| |
| self.paired_head_layers = [0, 2, 5, 9] |
| self.attn = CausalSelfAttention(model_dim, head_dim, num_heads, paired=False) |
| self.attn_paired = CausalSelfAttention(model_dim, head_dim, num_heads, paired=True) |
| self.yarn = Yarn(head_dim, max_seq_len) |
| self.yarn_paired_head = Yarn(head_dim, max_seq_len, paired=True) |
|
|
| |
| |
| |
| self.value_embeds = nn.Parameter(0.01 * torch.randn(5 * self.vocab_size, model_dim, dtype=torch.bfloat16)) |
|
|
| |
| self.attn_gate_bank = nn.Parameter(torch.zeros(10, num_heads, 12)) |
| self.ve_gate_bank = nn.Parameter(torch.zeros(5, num_heads, 12)) |
| self.gate_filler_nones = [None] * (num_layers - 6) |
|
|
| |
| |
| num_attn_layers = num_layers - 1 |
| hdim = num_heads * head_dim |
|
|
| |
| |
| self._num_attn_layers = num_attn_layers |
| num_qk_groups = num_attn_layers * 2 * (num_heads // 2) |
| self._num_qk_groups = num_qk_groups |
| num_qk_padded = next_multiple_of_n(num_qk_groups, n=world_size) |
| self.qk_bank = nn.Parameter(torch.empty(num_qk_padded, head_dim * 2, model_dim)) |
| self.qk_bank.reshape = (num_qk_padded, head_dim * 2, model_dim) |
|
|
| |
| num_vo_real = num_attn_layers * 2 |
| num_vo_padded = next_multiple_of_n(num_vo_real, n=world_size) |
| self.vo_bank = nn.Parameter(torch.empty(num_vo_padded, hdim, hdim)) |
| self.vo_bank.reshape = (num_vo_padded, hdim, hdim) |
|
|
| |
| std = 0.5 * model_dim ** -0.5 |
| bound = (3 ** 0.5) * std |
| with torch.no_grad(): |
| self.qk_bank[:num_qk_groups].uniform_(-bound, bound) |
| self.qk_bank[num_qk_groups:].zero_() |
| self.vo_bank[:num_vo_real].uniform_(-bound, bound) |
| self.vo_bank[num_vo_real:].zero_() |
|
|
| def init_mlp(self, model_dim): |
| |
| |
| mlp_hdim = 4 * model_dim |
| self.mlp_bank = nn.Parameter(torch.empty(12, 2, mlp_hdim, model_dim)) |
| self.mlp_bank.reshape = (24, mlp_hdim, model_dim) |
|
|
| |
| std = 0.5 * model_dim ** -0.5 |
| bound = (3 ** 0.5) * std |
| with torch.no_grad(): |
| self.mlp_bank[:, 0, :, :].uniform_(-bound, bound) |
| self.mlp_bank[:, 1, :, :].zero_() |
|
|
| def init_misc(self, model_dim, num_layers): |
| self.smear_gate = nn.Linear(12, 1, bias=False) |
| nn.init.zeros_(self.smear_gate.weight) |
|
|
| self.skip_gate = nn.Linear(12, 1, bias=False) |
| nn.init.zeros_(self.skip_gate.weight) |
|
|
| self.bigram_embed = nn.Embedding(args.bigram_vocab_size, args.bigram_dim) |
| nn.init.zeros_(self.bigram_embed.weight) |
| bigram_sign_table = torch.randn(args.bigram_sign_table_rows, args.bigram_dim).sign().to(torch.bfloat16) |
| self.register_buffer('bigram_sign_table', bigram_sign_table) |
|
|
| self.post_lambdas = nn.Parameter(torch.ones(num_layers, 2)) |
|
|
| |
| self.x0_lambdas = nn.Parameter(torch.zeros(num_layers)) |
| self.bigram_lambdas = nn.Parameter(0.05 * torch.ones(num_layers)) |
|
|
| |
| |
| self.resid_lambdas = nn.Parameter(torch.full((num_layers, 2), 1.1**0.5)) |
|
|
| |
| self.xsa_alphas = nn.Parameter(torch.zeros(num_layers, self.num_heads)) |
|
|
| pad = (-num_layers * 2 - 2) % dist.get_world_size() |
| self.scalars = nn.Parameter( |
| torch.cat( |
| [ |
| *[torch.tensor([0.5, 1.0]) for _ in range(num_layers)], |
| torch.zeros(1), |
| -1.5 * torch.ones(1), |
| torch.ones(pad), |
| ] |
| ) |
| ) |
|
|
| def init_mudd(self, num_layers: int, model_dim: int): |
| """ |
| Multiway Dynamic Dense Connections @lishengping. https://arxiv.org/abs/2502.12170 |
| Expressive and efficient mechanism for data dependent skip connections. |
| Given current activation x, return n skip coefficients computed via ~mlp(x). |
| Trimmed for speedrun: invoked at start of last layer and post-loop only. |
| |
| Start of last layer produces 14 coefficients: |
| mu[0..2] = v_mudd source coefs (cache[0], cache[7], x) -> added into V |
| mu[3..5] = residual source coefs (cache[0], cache[7], x) -> residual recombination |
| mu[6..7] = per-pair ve_gate (2 channels, tiled to num_heads) |
| mu[8..9] = resid_attn / post_attn lambdas (dynamic) |
| mu[10..11]= x0 / bigram injection lambdas (dynamic) |
| mu[12..13]= resid_mlp / post_mlp lambdas (dynamic) |
| |
| Post-loop produces 5 residual coefs over |
| {cache[0], cache[7], cache[9], ve_bank0, cache[3]}. |
| """ |
| num_mudd_layers = 2 |
| self._mudd_scale = 0.1 |
| mudd_dim = 64 |
| max_num_coef = 14 |
|
|
| self.mudd_w1 = nn.Parameter(torch.empty(num_mudd_layers, mudd_dim, model_dim)) |
| for j in range(num_mudd_layers): |
| nn.init.kaiming_uniform_(self.mudd_w1.data[j], a=math.sqrt(5)) |
|
|
| self.mudd_w2 = nn.Parameter(torch.zeros(num_mudd_layers, max_num_coef, mudd_dim)) |
|
|
| |
| bs_init = torch.zeros(num_mudd_layers, max_num_coef) |
| |
| bs_init[0, 6] = 2.0 / self._mudd_scale |
| bs_init[0, 7] = 2.0 / self._mudd_scale |
| |
| bs_init[0, 8] = 1.1**0.5 / self._mudd_scale |
| bs_init[0, 9] = 1.0 / self._mudd_scale |
| bs_init[0, 10] = 0.0 |
| bs_init[0, 11] = 0.05 / self._mudd_scale |
| bs_init[0, 12] = 1.1**0.5 / self._mudd_scale |
| bs_init[0, 13] = 1.0 / self._mudd_scale |
| |
| bs_init[1, 1] = -0.5 / self._mudd_scale |
| self.mudd_b2 = nn.Parameter(bs_init) |
|
|
| def forward_mudd(self, x, id, num_coef): |
| """Returns `num_coef` per-token MUDD coefficients from block `id` (0 or 1).""" |
| x = F.gelu(F.linear(x, self.mudd_w1[id])) |
| x = (F.linear(x, self.mudd_w2[id, :num_coef]) + self.mudd_b2[id, :num_coef]) * self._mudd_scale |
| return x.split(1, dim=-1) |
|
|
| def forward(self, input_seq: Tensor, target_seq: Tensor, seqlens: Tensor, bigram_input_seq: Tensor, schedule_cfg: ForwardScheduleConfig): |
| assert input_seq.ndim == 1 |
|
|
| |
| mtp_weights, train_max_seq_len = schedule_cfg.mtp_weights, schedule_cfg.train_max_seq_len |
| ws_short, ws_long = schedule_cfg.ws_short, schedule_cfg.ws_long |
| |
| bm_sizes = [ws_short, ws_short, ws_short, ws_long, ws_short, ws_short, None, ws_short, ws_short, ws_short, ws_long] |
| assert len(bm_sizes) == self.num_layers |
| key_offset = [b==ws_long for b in bm_sizes] |
|
|
| |
| sa_lambdas = self.scalars[: 2 * self.num_layers].view(-1, 2) |
| smear_lambda = self.scalars[2 * self.num_layers] |
| skip_lambda = self.scalars[2 * self.num_layers + 1] |
| resid_lambdas_attn = self.resid_lambdas[:, 0].bfloat16().unbind(0) |
| resid_lambdas_mlp = self.resid_lambdas[:, 1].bfloat16().unbind(0) |
| post_lambdas_attn = self.post_lambdas[:, 0].bfloat16().unbind(0) |
| post_lambdas_mlp = self.post_lambdas[:, 1].bfloat16().unbind(0) |
| x0_lambdas = self.x0_lambdas.bfloat16().unbind(0) |
| bigram_lambdas = self.bigram_lambdas.bfloat16().unbind(0) |
| ag = self.attn_gate_bank.unbind(0) |
| veg = self.ve_gate_bank.unbind(0) |
| attn_gates = [*ag[:6], None, *ag[6:]] |
| ve_gates = [None, veg[0], veg[1], *self.gate_filler_nones, veg[2], veg[3], veg[4]] |
| |
| xsa_alpha_per_layer = self.xsa_alphas.unbind(0) |
| xsa_alphas = [xsa_alpha_per_layer[j] if j in {1, 3, 4, 7, 8, 10} else None for j in range(self.num_layers)] |
| assert len(attn_gates) == self.num_layers |
| assert len(ve_gates) == self.num_layers |
| qk_all = self.qk_bank[:self._num_qk_groups].view(self._num_attn_layers, -1, self.qk_bank.shape[-1]) |
| vo_flat = self.vo_bank[:self._num_attn_layers * 2].view(self._num_attn_layers, 2, *self.vo_bank.shape[1:]).flatten(1, 2) |
| attn_weights = torch.cat([qk_all, vo_flat], dim=1).unbind(0) |
| mlp_all = self.mlp_bank.flatten(0, 1).unbind(0) |
| mlp_fcs = mlp_all[0::2] |
| mlp_projs = mlp_all[1::2] |
|
|
| |
| x = self.embed(input_seq) |
| |
| |
| |
| sign_idx = torch.zeros_like(input_seq) |
| sign_idx[1:] = (input_seq[:-1] ^ input_seq[1:]) % self.bigram_sign_table.shape[0] |
| bigram_signs = self.bigram_sign_table[sign_idx] |
| x0_bigram = (self.bigram_embed(bigram_input_seq) * bigram_signs)[None] |
|
|
| |
| ve = self.value_embeds.view(5, self.vocab_size, -1)[:, input_seq] |
| |
| ve = [None, ve[0], ve[1], *self.gate_filler_nones, ve[2], ve[3], ve[4]] |
| assert len(ve) == self.num_layers |
|
|
| |
| smear_gate_out = smear_lambda * torch.sigmoid(self.smear_gate(x[1:, :self.smear_gate.weight.size(-1)])) |
| x = torch.cat([x[:1], x[1:] + smear_gate_out * x[:-1]]) |
| x = x0 = norm(x[None]) |
|
|
| |
| x[..., :args.bigram_dim] = x[..., :args.bigram_dim] + x0_bigram * bigram_lambdas[0] |
|
|
| |
| |
| x0_inject = tuple(x0 * x0_lambdas[i] for i in range(self.num_layers)) |
| bg_inject = (None,) + tuple(x0_bigram * bigram_lambdas[i] for i in range(1, self.num_layers)) |
| skip_gate_out = torch.sigmoid(skip_lambda) * 2 * torch.sigmoid(self.skip_gate(x0[..., :self.skip_gate.weight.size(-1)])) |
|
|
| |
| |
| cache = {0: x} |
| for i in range(self.num_layers): |
| is_paired = i in self.paired_head_layers |
| yarn = self.yarn_paired_head if is_paired else self.yarn |
| attn = self.attn_paired if is_paired else self.attn |
| c_fc = mlp_fcs[i] |
| c_proj = mlp_projs[i] |
| mu = None |
|
|
| |
| if i == 6: |
| x = x + skip_gate_out * cache[3] |
| else: |
| qkvo_w = attn_weights[i - (i > 6)] |
| attn_in_normed = norm(cache.get(7, x)) |
| B, T = attn_in_normed.size(0), attn_in_normed.size(1) |
|
|
| if i == self.num_layers - 1: |
| cache[9] = x |
| mu = self.forward_mudd(x, id=0, num_coef=14) |
| v_mudd = (mu[0] * cache[0] + mu[1] * cache[7] + mu[2] * x).view(B, T, self.num_heads, self.head_dim) |
| x = (1 + mu[5]) * x + mu[3] * cache[0] + mu[4] * cache[7] |
| ve_gate = torch.cat([mu[6], mu[7]], dim=-1).repeat_interleave( |
| self.num_heads // 2, dim=-1 |
| ).unsqueeze(-1) |
| ve_view = ve[i].view(B, T, self.num_heads, self.head_dim) |
| aux_v = (ve_gate * ve_view + v_mudd).view(B, T, -1) |
| elif ve[i] is not None: |
| |
| gate_in = torch.cat([attn_in_normed[..., :6], ve[i][None, ..., :6]], dim=-1) |
| ve_gate_out = 2 * torch.sigmoid(F.linear(gate_in, ve_gates[i])).view(B, T, self.num_heads, 1) |
| ve_view = ve[i].view(B, T, self.num_heads, self.head_dim) |
| aux_v = (ve_gate_out * ve_view).view(B, T, -1) |
| else: |
| aux_v = None |
|
|
| attn_args = AttnArgs( |
| sa_lambdas=sa_lambdas[i], |
| seqlens=seqlens, |
| bm_size=bm_sizes[i], |
| yarn=yarn, |
| key_offset=key_offset[i], |
| attn_gate_w=attn_gates[i], |
| aux_v=aux_v, |
| xsa_alpha=xsa_alphas[i], |
| train_max_seq_len=train_max_seq_len, |
| ) |
| attn_out = attn(attn_in_normed, attn_args, qkvo_w) |
|
|
| if mu is not None: |
| x = mu[8] * x + mu[9] * attn_out + mu[10] * cache[0] |
| x[..., :args.bigram_dim] = x[..., :args.bigram_dim] + mu[11] * x0_bigram |
| else: |
| x = resid_lambdas_attn[i] * x + post_lambdas_attn[i] * attn_out + x0_inject[i] |
| if bg_inject[i] is not None: |
| x[..., :args.bigram_dim] = x[..., :args.bigram_dim] + bg_inject[i] |
|
|
| if mu is not None: |
| x = mu[12] * x + mu[13] * ReLUSqrdMLP(norm(x), c_fc, c_proj) |
| else: |
| x = resid_lambdas_mlp[i] * x + post_lambdas_mlp[i] * ReLUSqrdMLP(norm(x), c_fc, c_proj) |
|
|
| if i in self.cache_layers: |
| cache[i] = x |
|
|
| |
| mu = self.forward_mudd(x, id=1, num_coef=5) |
| ve_bank0 = ve[1][None].to(dtype=x.dtype) |
| x = x + mu[0] * cache[0] + mu[1] * cache[7] + mu[2] * cache[9] + mu[3] * ve_bank0 + mu[4] * cache[3] |
|
|
| x = norm(x) |
| |
| |
| if self.training: |
| loss_per_token = FusedSoftcappedCrossEntropy.apply(x.view(-1, x.size(-1)), target_seq, mtp_weights, self.lm_head.weight, self.lm_head.x_s, self.lm_head.w_s, self.lm_head.grad_s, grad_scale) |
| else: |
| logits = self.lm_head(x) |
| logits = 23 * torch.sigmoid((logits + 5) / 7.5) |
| logits_for_loss = logits.float() |
| loss_per_token = F.cross_entropy(logits_for_loss.view(-1, logits_for_loss.size(-1)), target_seq, reduction="none") |
| return loss_per_token |
| |
| |
|
|
| def _load_data_shard(file: Path): |
| header = torch.from_file(str(file), False, 256, dtype=torch.int32) |
| assert header[0] == 20240520, "magic number mismatch in the data .bin file" |
| assert header[1] == 1, "unsupported version" |
| num_tokens = int(header[2]) |
| with file.open("rb", buffering=0) as f: |
| tokens = torch.empty(num_tokens, dtype=torch.uint16, pin_memory=True) |
| f.seek(256 * 4) |
| nbytes = f.readinto(tokens.numpy()) |
| assert nbytes == 2 * num_tokens, "number of tokens read does not match header" |
| return tokens |
|
|
| BOS_ID = 0 |
| TRAIN_MAX_NUM_DOCS = {16384: 384, 32768: 768, 49152: 1152} |
|
|
| class Shard: |
| def __init__(self, tokens: Tensor, world_size: int = 1): |
| self.tokens = tokens |
| self.size = tokens.numel() |
| self.world_size = world_size |
| self.i = 0 |
|
|
| |
| self.bos_idx = (tokens[:6_000_000] == BOS_ID).nonzero(as_tuple=True)[0].to(torch.int64).cpu().numpy() |
| self._full_idx = None |
| self._loader_thread = None |
| self._ready = threading.Event() |
| self._loader_thread = threading.Thread(target=self._scan) |
| self._loader_thread.start() |
|
|
| def _scan(self): |
| self._full_idx = (self.tokens == BOS_ID).nonzero(as_tuple=True)[0].to(torch.int64).cpu().numpy() |
| self._ready.set() |
|
|
| def _maybe_switch(self): |
| |
| if self.bos_idx is not self._full_idx and self._ready.is_set(): |
| self._loader_thread.join() |
| self.bos_idx = self._full_idx |
|
|
| def next_batch(self, num_tokens_local: int, max_seq_len: int): |
| self._maybe_switch() |
| n = len(self.bos_idx) |
| starts = [[] for _ in range(self.world_size)] |
| ends = [[] for _ in range(self.world_size)] |
|
|
| idx = self.i |
| for r in range(self.world_size): |
| cur_len = 0 |
| while cur_len <= num_tokens_local: |
| if idx >= n: |
| raise StopIteration(f"Insufficient BOS ahead; hit tail of shard.") |
| cur = self.bos_idx[idx] |
| starts[r].append(cur) |
| idx += 1 |
| end = min(self.bos_idx[idx] if idx < n else self.size, |
| cur + max_seq_len, |
| cur + num_tokens_local - cur_len + 1) |
| ends[r].append(end) |
| cur_len += end - cur |
|
|
| assert cur_len == num_tokens_local + 1 |
| self.i = idx |
| return starts, ends |
|
|
| @staticmethod |
| def load_async(file: Path, world_size: int = 1): |
| """Returns getter function for async shard loading""" |
| result = {} |
| ready = threading.Event() |
| def load(): |
| tokens = _load_data_shard(file) |
| result['shard'] = Shard(tokens, world_size) |
| ready.set() |
| thread = threading.Thread(target=load) |
| thread.start() |
| def get(): |
| ready.wait() |
| thread.join() |
| return result['shard'] |
| return get |
|
|
| def get_bigram_hash(x): |
| """ |
| Computes bigram hash for each position using [prev_token, curr_token]. |
| Multiply by arbitary large ints to get even spread over int32 range. |
| Position 0 is mapped to the reserved index (vocab_size - 1). |
| BOS_tokens within the batch will hash based on last token of prior doc. Masking this ran slower and showed no improvement. |
| """ |
| rand_int_1 = 36313 |
| rand_int_2 = 27191 |
| mod = args.bigram_vocab_size-1 |
| x = x.to(torch.int32) |
| out = torch.empty_like(x, pin_memory=True) |
| out.copy_(x) |
| out[0] = mod |
| out[1:] = torch.bitwise_xor(rand_int_1 * out[1:], rand_int_2 * out[:-1]) % mod |
| return out |
|
|
| def distributed_data_generator(filename_pattern: str, num_tokens: int, max_seq_len: int, grad_accum_steps: int = 1, align_to_bos: bool = True): |
| |
| rank = dist.get_rank() if dist.is_initialized() else 0 |
| world_size = dist.get_world_size() if dist.is_initialized() else 1 |
| assert num_tokens % (world_size * grad_accum_steps) == 0, "Batch size must be divisible by world size" |
| num_tokens = num_tokens // grad_accum_steps |
|
|
| files = [Path(file) for file in sorted(glob.glob(filename_pattern))] |
| if not files: |
| raise FileNotFoundError(f"No files found for pattern: {filename_pattern}") |
|
|
| file_iter = iter(files) |
| tokens = _load_data_shard(next(file_iter)) |
| if align_to_bos: |
| shard = Shard(tokens, world_size) |
| next_shard_getter = Shard.load_async(next(file_iter), world_size) |
| else: |
| pos = 0 |
|
|
| while True: |
| num_tokens_local = num_tokens // world_size |
| max_num_docs = TRAIN_MAX_NUM_DOCS.get(num_tokens_local, next_multiple_of_n(num_tokens_local // 48, n=128)) |
|
|
| if align_to_bos: |
| try: |
| seq_starts, seq_ends = shard.next_batch(num_tokens_local, max_seq_len) |
| start_idxs, end_idxs = torch.tensor(seq_starts[rank]), torch.tensor(seq_ends[rank]) |
| except StopIteration: |
| |
| shard = next_shard_getter() |
| tokens = shard.tokens |
| try: |
| next_shard_getter = Shard.load_async(next(file_iter), world_size) |
| except StopIteration: |
| next_shard_getter = None |
| continue |
|
|
| buf = torch.cat([tokens[i:j] for i, j in zip(start_idxs, end_idxs)]) |
| _inputs = buf[:-1] |
| _targets = buf[1:] |
| end_idxs[-1] -= 1 |
| cum_lengths = (end_idxs - start_idxs).cumsum(0) |
|
|
| else: |
| if pos + num_tokens + 1 >= len(tokens): |
| tokens, pos = _load_data_shard(next(file_iter)), 0 |
|
|
| pos_local = pos + rank * num_tokens_local |
| buf = tokens[pos_local: pos_local + num_tokens_local + 1] |
| _inputs = buf[:-1].view(num_tokens_local, ) |
| _targets = buf[1:].view(num_tokens_local, ) |
|
|
| cum_lengths = torch.nonzero(_inputs == BOS_ID)[:, 0] |
| pos += num_tokens |
|
|
|
|
| _cum_lengths = torch.full((max_num_docs,), num_tokens_local) |
| _cum_lengths[0] = 0 |
| _cum_lengths[1:len(cum_lengths) + 1] = cum_lengths |
|
|
| |
| _inputs = _inputs.to(dtype=torch.int32) |
| _targets = _targets.to(dtype=torch.int64) |
| _cum_lengths = _cum_lengths.to(dtype=torch.int32) |
| _bigram_inputs = get_bigram_hash(_inputs) |
|
|
| new_params = yield ( |
| _inputs.to(device="cuda", non_blocking=True), |
| _targets.to(device="cuda", non_blocking=True), |
| _cum_lengths.to(device="cuda", non_blocking=True), |
| _bigram_inputs.to(device="cuda", non_blocking=True), |
| _bigram_inputs.numpy(), |
| ) |
|
|
| if new_params is not None: |
| |
| new_num_tokens, new_max_seq_len, new_grad_accum_steps = new_params |
| assert new_num_tokens % (world_size * new_grad_accum_steps) == 0, "Num tokens must be divisible by world size" |
| num_tokens = new_num_tokens // new_grad_accum_steps |
| max_seq_len = new_max_seq_len |
|
|
| |
| |
|
|
| @dataclass(slots=True) |
| class Hyperparameters: |
| |
| data_path = os.environ.get("DATA_PATH", ".") |
| train_files: str = os.path.expanduser("~/dynaword/shards/polish_train_*.bin") |
| val_files: str = os.path.expanduser("~/dynaword/shards/polish_val_*.bin") |
| val_tokens: int = 10485760 |
| |
| val_batch_size: int = 4 * 64 * 1024 * 8 |
| |
| num_scheduled_iterations: int = 13200 |
| num_extension_iterations: int = 10 |
| |
| run_id: str = f"{uuid.uuid4()}" |
| |
| |
| |
| |
| val_loss_every: int = 250 |
| save_checkpoint: bool = True |
| checkpoint_every: int = 500 |
| run_evals: bool = False |
| |
| bigram_vocab_size: int = 50304 * 15 |
| bigram_dim: int = 192 |
| bigram_sign_table_rows: int = 8192 |
|
|
| args = Hyperparameters() |
|
|
| @dataclass(slots=True) |
| class TrainingStage: |
| lr_mul: float |
| batch_size: int |
| window_sizes: tuple[int, int] |
| mtp_weights_start: list[float] |
| mtp_weights_end: list[float] |
| train_max_seq_len: int |
| duration: float = None |
|
|
| class TrainingSchedule: |
| """ |
| Training schedule initialized via TRAINING_STAGES |
| 1. Multi Token Prediction schedule of [1, 0.5, 0.25->0] -> [1, 0.5->0] -> [1] @varunneal |
| 2. Sliding Attention window schedule of [1,3] -> [3,7] -> [5,11] -> [6,13] |
| 3. YaRN updates to RoPE on window changes |
| 4. Split embed and lm head at 2/3 of training |
| 5. Batch size schedule of 8 -> 16 -> 24 |
| 6. Post training extension of long windows from 13 to 20 |
| 7. Seq len updates from 896 to 2048 at 1/3 of training |
| """ |
|
|
| def __init__(self, stages: list[TrainingStage], scheduled_iterations: int, extension_iterations: int, |
| cooldown_frac: float = 0.5, split_embed_stage: int = 2, ws_post_yarn_ext: int = 20): |
| self.stages = stages |
| self.scheduled_iterations = scheduled_iterations |
| self.cooldown_frac = cooldown_frac |
| |
| self.ws_post_yarn_ext = ws_post_yarn_ext |
|
|
| self.total_steps = self.scheduled_iterations + extension_iterations |
|
|
| |
| ends = [0, *[round(c * scheduled_iterations) for c in accumulate(s.duration for s in stages[:-1])], self.total_steps] |
| assert self.scheduled_iterations == ends[-2] |
| self.boundaries = list(pairwise(ends)) |
|
|
| |
| self.split_step = self.boundaries[split_embed_stage][0] | 1 |
|
|
| |
| self.mtp_weights = [] |
| for step in range(self.total_steps + 1): |
| stage, t = self.lookup(step) |
| w = [a + (b - a) * t for a, b in zip(stage.mtp_weights_start, stage.mtp_weights_end)] |
| self.mtp_weights.append(torch.tensor(w, device=device)) |
|
|
| def lookup(self, step: int) -> tuple[TrainingStage, float]: |
| |
| for i, (start, end) in enumerate(self.boundaries): |
| if step < end: |
| t = (step - start) / (end - start) |
| return self.stages[i], t |
| return self.stages[-1], 1.0 |
|
|
| def get_lr(self, step: int) -> float: |
| |
| stage, _ = self.lookup(step) |
| lr = stage.lr_mul |
| cd_start = int(self.scheduled_iterations * (1 - self.cooldown_frac)) |
| if step >= cd_start: |
| t = min(1.0, (step - cd_start) / (self.scheduled_iterations - cd_start)) |
| lr = lr * (1 - t) + 0.15 * t |
| return lr |
|
|
| |
| TRAINING_STAGES = [ |
| TrainingStage(duration=1/3, train_max_seq_len=896, batch_size=8 * 2048 * 8, window_sizes=(1, 3), lr_mul=1.0, |
| mtp_weights_start=[1.0, 0.5, 0.25], mtp_weights_end=[1.0, 0.5, 0.0]), |
| TrainingStage(duration=1/3, train_max_seq_len=2048, batch_size=16 * 2048 * 8, window_sizes=(3, 7), lr_mul=1.52, |
| mtp_weights_start=[1.0, 0.5], mtp_weights_end=[1.0, 0.0]), |
| TrainingStage(duration=1/3, train_max_seq_len=2048, batch_size=24 * 2048 * 8, window_sizes=(5, 11), lr_mul=1.73, |
| mtp_weights_start=[1.0], mtp_weights_end=[1.0]), |
| |
| TrainingStage(train_max_seq_len=2048, batch_size=24 * 2048 * 8, window_sizes=(6, 13), lr_mul=1.0, |
| mtp_weights_start=[1.0], mtp_weights_end=[1.0]), |
| ] |
|
|
| |
| training_schedule = TrainingSchedule(TRAINING_STAGES, args.num_scheduled_iterations, args.num_extension_iterations, cooldown_frac=0.60) |
| |
|
|
| def get_muon_momentum(step: int, muon_warmup_steps=300, muon_cooldown_steps=50, momentum_min=0.85, momentum_max=0.95): |
| |
| |
| momentum_cd_start = training_schedule.total_steps - muon_cooldown_steps |
| if step < muon_warmup_steps: |
| frac = step / muon_warmup_steps |
| momentum = momentum_min + frac * (momentum_max - momentum_min) |
| elif step > momentum_cd_start: |
| frac = (step - momentum_cd_start) / muon_cooldown_steps |
| momentum = momentum_max - frac * (momentum_max - momentum_min) |
| else: |
| momentum = momentum_max |
| return momentum |
|
|
| class TrainingManager(): |
| """ |
| Manages the NorMuonAndAdam for all parameters with explicit ordering. |
| 1. Scalars are given higher momentum terms to smooth learning @ChrisJMcCormick |
| 2. Adam optimizers are only stepped on odd steps @classiclarryd |
| 3. Explicit scatter_order and work_order for communication scheduling (no backward hooks) |
| 4. Muon has a linear momentum warmup and cooldown schedule |
| 5. Learning rates follow a linear decay schedule |
| 6. Embed is tied to lm_head until split step (2/3 of training), then untied @classiclarryd |
| """ |
| def __init__(self, model): |
| self.model = model |
| self.block_size = 128 |
|
|
| |
| |
| |
| self.param_table = { |
| "qk_bank": {"optim": "normuon", "comms": "sharded", "adam_betas": None}, |
| "vo_bank": {"optim": "normuon", "comms": "sharded", "adam_betas": None}, |
| "mlp_bank": {"optim": "normuon", "comms": "sharded", "adam_betas": None}, |
| "scalars": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 5.0, "wd_mul": 0.0}, |
| "smear_gate": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 0.01, "wd_mul": 0.0}, |
| "skip_gate": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 0.05, "wd_mul": 0.0}, |
| "attn_gate_bank": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99]}, |
| "ve_gate_bank": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99]}, |
| "lm_head": {"optim": "adam", "comms": "sharded", "adam_betas": [0.5, 0.95], "wd_mul": 150.}, |
| "bigram_embed": {"optim": "adam", "comms": "sharded_sparse", "adam_betas": [0.75, 0.95], "lr_mul": 75., "wd_mul": 5.0}, |
| "post_lambdas": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.95], "lr_mul": 1.0, "wd_mul": 0.0}, |
| "x0_lambdas": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.95], "lr_mul": 1.0, "wd_mul": 0.0}, |
| "bigram_lambdas": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.95], "lr_mul": 1.0, "wd_mul": 0.0}, |
| "resid_lambdas": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.95], "lr_mul": 5.0, "wd_mul": 0.0}, |
| "xsa_alphas": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.95], "lr_mul": 1.0, "wd_mul": 0.0}, |
| "value_embeds": {"optim": "adam", "comms": "sharded", "adam_betas": [0.75, 0.95], "lr_mul": 75., "wd_mul": 5.0}, |
| "embed": {"optim": "adam", "comms": "sharded", "adam_betas": [0.5, 0.95], "wd_mul": 150.}, |
| } |
|
|
| |
| self.param_table.update({ |
| "mudd_w1": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 0.25}, |
| "mudd_w2": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 0.25}, |
| "mudd_b2": {"optim": "adam", "comms": "replicated", "adam_betas": [0.9, 0.99], "lr_mul": 0.25, "wd_mul": 0.0}, |
| }) |
|
|
| |
| |
| self.work_order = [ |
| "scalars", "smear_gate", "skip_gate", "attn_gate_bank", "ve_gate_bank", "mudd_b2", "xsa_alphas", |
| "post_lambdas", "x0_lambdas", "bigram_lambdas", "resid_lambdas", |
| ] + [ |
| "mudd_w2", |
| "value_embeds", "bigram_embed", |
| "mudd_w1", |
| "lm_head", "embed", |
| "qk_bank", "vo_bank", "mlp_bank", |
| ] |
|
|
| adam_defaults = dict( |
| lr=0.008, |
| eps=1e-10, |
| weight_decay=0.005, |
| ) |
|
|
| normuon_defaults = dict( |
| lr=0.023, |
| momentum=0.95, |
| beta2=0.9, |
| weight_decay=1.2, |
| ) |
|
|
| self.optimizer = NorMuonAndAdam( |
| model.named_parameters(), |
| param_table=self.param_table, |
| scatter_order=list(self.param_table), |
| work_order=self.work_order, |
| adam_defaults=adam_defaults, |
| normuon_defaults=normuon_defaults, |
| ) |
|
|
| |
| self.split_step = training_schedule.split_step |
|
|
| self.reset() |
|
|
| def apply_final_ws_ext(self): |
| self.ws_long = training_schedule.ws_post_yarn_ext |
|
|
| def get_forward_args(self): |
| return ForwardScheduleConfig( |
| mtp_weights = self.mtp_weights, |
| ws_short = self.ws_short * self.block_size, |
| ws_long = self.ws_long * self.block_size, |
| train_max_seq_len = self.train_max_seq_len |
| ) |
|
|
| def _is_adam_step(self, step: int): |
| """Adam params are only updated on odd steps.""" |
| return step % 2 == 1 |
|
|
| def get_transition_steps(self): |
| return [start for start, _ in training_schedule.boundaries[1:]] |
|
|
| def advance_schedule(self, step: int): |
| stage, _ = training_schedule.lookup(step) |
| self.ws_short, new_ws_long = stage.window_sizes |
| if new_ws_long != self.ws_long: |
| self.model.yarn.apply(self.ws_long * self.block_size, new_ws_long * self.block_size) |
| self.model.yarn_paired_head.apply(self.ws_long * self.block_size, new_ws_long * self.block_size) |
|
|
| new_batch_size = stage.batch_size |
| new_train_max_seq_len = stage.train_max_seq_len |
| if new_batch_size != self.batch_size or new_train_max_seq_len != self.train_max_seq_len: |
| self.train_loader_send_args = (new_batch_size, new_train_max_seq_len, grad_accum_steps) |
| self.batch_size = new_batch_size |
| self.train_max_seq_len = new_train_max_seq_len |
| else: |
| self.train_loader_send_args = None |
|
|
| self.ws_long = new_ws_long |
| self.mtp_weights = training_schedule.mtp_weights[step] |
|
|
| def step_optimizers(self, step: int): |
| step_lr = training_schedule.get_lr(step) |
| muon_momentum = get_muon_momentum(step) |
| do_adam = self._is_adam_step(step) |
|
|
| |
| for param, p_cfg in self.optimizer.param_cfgs.items(): |
| p_cfg.lr = p_cfg.initial_lr * step_lr |
| if p_cfg.optim == "normuon": |
| p_cfg.momentum = muon_momentum |
|
|
| |
| self.optimizer.step(do_adam=do_adam) |
|
|
| |
| if step == self.split_step: |
| self.optimizer.copy_lm_state_to_embed() |
|
|
| def reset(self, state=None): |
| if state is not None: |
| self.optimizer.load_state_dict(state) |
|
|
| |
| self.optimizer.reset() |
|
|
| stage, _ = training_schedule.lookup(0) |
| self.ws_short, self.ws_long = stage.window_sizes |
| self.batch_size = stage.batch_size |
| self.train_max_seq_len = stage.train_max_seq_len |
| self.model.yarn.reset() |
| self.model.yarn_paired_head.reset() |
| if _sparse_comms_active(): |
| self.row_update_mask = np.zeros(args.bigram_vocab_size, dtype=np.uint8) |
| self.sparse_counts_state = None |
| |
| self.send_idxes_buffer = torch.empty(args.bigram_vocab_size, dtype=torch.int32, pin_memory=True) |
|
|
|
|
| def get_state(self): |
| return copy.deepcopy(self.optimizer.state_dict()) |
|
|
| def sparse_index_update(self, step, bigram_indexes): |
| if not _sparse_comms_active(): |
| return |
|
|
| self.row_update_mask[bigram_indexes] = 1 |
|
|
| if self._is_adam_step(step): |
| with torch.no_grad(): |
| bigram_idx_np = np.flatnonzero(self.row_update_mask).astype(np.int32) |
| send_idxes, send_counts, recv_counts, recv_counts_fut = sparse_comms_start( |
| bigram_idx_np, args.bigram_vocab_size, rank, world_size, self.send_idxes_buffer |
| ) |
| self.sparse_counts_state = (send_idxes, send_counts, recv_counts, recv_counts_fut) |
|
|
| def sparse_index_share(self, step): |
| if not _sparse_comms_active() or not self._is_adam_step(step): |
| return |
|
|
| send_idxes, send_counts, recv_counts, recv_counts_fut = self.sparse_counts_state |
| self.sparse_counts_state = None |
|
|
| recv_counts_fut.wait() |
| recv_idxes, sparse_state, idxes_fut = sparse_comms_share_indexes(send_idxes, send_counts, recv_counts) |
| self.optimizer._reduce_futures[model.bigram_embed.weight] = [idxes_fut, recv_idxes] |
| self.optimizer._sparse_async_data[model.bigram_embed.weight] = sparse_state |
|
|
| self.row_update_mask.fill(0) |
|
|
|
|
| |
|
|
| |
| |
|
|
| |
| logfile = None |
| if master_process: |
| run_id = args.run_id |
| os.makedirs("logs", exist_ok=True) |
| logfile = f"logs/{run_id}.txt" |
| print(logfile) |
| def print0(s, console=False): |
| if master_process: |
| with open(logfile, "a") as f: |
| if console: |
| print(s) |
| print(s, file=f) |
|
|
| |
| print0(code) |
| print0("="*100) |
| |
| print0(f"Running Python {sys.version}") |
| print0(f"Running PyTorch {torch.version.__version__} compiled for CUDA {torch.version.cuda}") |
| print0(f"Running Triton version {triton.__version__}") |
|
|
| def nvidia_smi(): |
| import subprocess |
| return subprocess.run(["nvidia-smi"], stdout=subprocess.PIPE, stderr=subprocess.PIPE, text=True).stdout |
| print0(nvidia_smi()) |
| print0("="*100) |
|
|
| model: nn.Module = GPT( |
| vocab_size=32896, |
| num_layers=11, |
| num_heads=6, |
| head_dim=128, |
| model_dim=768, |
| max_seq_len=args.val_batch_size // (grad_accum_steps * world_size) |
| ).cuda() |
| for m in model.modules(): |
| if isinstance(m, (nn.Embedding, nn.Linear)): |
| m.weight.data = m.weight.data.bfloat16() |
| model.attn_gate_bank.data = model.attn_gate_bank.data.bfloat16() |
| model.ve_gate_bank.data = model.ve_gate_bank.data.bfloat16() |
| model.qk_bank.data = model.qk_bank.data.bfloat16() |
| model.vo_bank.data = model.vo_bank.data.bfloat16() |
| model.mlp_bank.data = model.mlp_bank.data.bfloat16() |
| model.mudd_w1.data = model.mudd_w1.data.bfloat16() |
| model.mudd_w2.data = model.mudd_w2.data.bfloat16() |
| model.mudd_b2.data = model.mudd_b2.data.bfloat16() |
| for param in model.parameters(): |
| dist.broadcast(param.detach(), 0) |
| dist.broadcast(model.bigram_sign_table, 0) |
|
|
| model: nn.Module = torch.compile(model, dynamic=False, fullgraph=True) |
| training_manager = TrainingManager(model) |
|
|
|
|
| |
| |
| |
| print0("Compiling model and warming up kernels (~7 minutes on first execution)", console=True) |
| |
| initial_state = dict(model=copy.deepcopy(model.state_dict()), |
| optimizer=training_manager.get_state()) |
| train_loader = distributed_data_generator(args.train_files, TRAINING_STAGES[0].batch_size, TRAINING_STAGES[0].train_max_seq_len, grad_accum_steps=grad_accum_steps) |
| val_loader = distributed_data_generator(args.val_files, args.val_batch_size, -1, grad_accum_steps=grad_accum_steps, align_to_bos=False) |
|
|
| transition_steps = training_manager.get_transition_steps() |
| |
| warmup_steps = sorted({0, 1} | {s + offset for s in transition_steps for offset in [-2, -1, 0, 1] if s + offset >= 2}) |
| print0(f"Sampling steps {warmup_steps} for warmup", console=True) |
| for step in warmup_steps: |
| training_manager.advance_schedule(step) |
| model.eval() |
| with torch.no_grad(): |
| inputs, targets, cum_seqlens, bigram_inputs, _ = next(val_loader) |
| model(inputs, targets, cum_seqlens, bigram_inputs, training_manager.get_forward_args()).mean() |
| model.train() |
| for idx in range(grad_accum_steps): |
| send_args = training_manager.train_loader_send_args |
| inputs, targets, cum_seqlens, bigram_inputs, bigram_cpu = train_loader.send(send_args) |
| training_manager.sparse_index_update(step, bigram_cpu) |
| loss = model(inputs, targets, cum_seqlens, bigram_inputs, training_manager.get_forward_args()).sum() * grad_scale |
| training_manager.sparse_index_share(step) |
| loss.backward() |
| del loss |
| training_manager.step_optimizers(step) |
| print0("Resetting Model", console=True) |
| model.zero_grad(set_to_none=True) |
| model.load_state_dict(initial_state["model"]) |
| training_manager.reset(initial_state["optimizer"]) |
| del val_loader, train_loader, initial_state |
| model.train() |
|
|
| |
| |
| |
| train_loader = distributed_data_generator(args.train_files, TRAINING_STAGES[0].batch_size, TRAINING_STAGES[0].train_max_seq_len, grad_accum_steps=grad_accum_steps) |
|
|
| gc.collect() |
|
|
| training_time_ms = 0 |
| |
| torch.cuda.synchronize() |
| t0 = time.perf_counter() |
| |
| train_steps = training_schedule.total_steps |
| for step in range(train_steps + 1): |
| last_step = (step == train_steps) |
| training_manager.advance_schedule(step) |
| |
| if last_step or (args.val_loss_every > 0 and step % args.val_loss_every == 0): |
| if last_step: |
| training_manager.apply_final_ws_ext() |
| |
| torch.cuda.synchronize() |
| training_time_ms += 1000 * (time.perf_counter() - t0) |
| model.eval() |
| assert args.val_tokens % args.val_batch_size == 0 |
| val_steps = grad_accum_steps * args.val_tokens // args.val_batch_size |
| val_loader = distributed_data_generator(args.val_files, args.val_batch_size, -1, grad_accum_steps=grad_accum_steps, align_to_bos=False) |
| val_loss = 0 |
| with torch.no_grad(): |
| for _ in range(val_steps): |
| inputs, targets, cum_seqlens, bigram_inputs, _ = next(val_loader) |
| val_loss += model(inputs, targets, cum_seqlens, bigram_inputs, training_manager.get_forward_args()).mean() |
| val_loss /= val_steps |
| del val_loader |
| dist.reduce(val_loss, 0, op=dist.ReduceOp.AVG) |
| print0(f"step:{step}/{train_steps} val_loss:{val_loss:.4f} train_time:{training_time_ms:.0f}ms step_avg:{training_time_ms/max(step, 1):.2f}ms", console=True) |
| model.train() |
| |
| torch.cuda.synchronize() |
| t0 = time.perf_counter() |
|
|
| if master_process and args.save_checkpoint and (last_step or (step > 0 and step % args.checkpoint_every == 0)): |
| log = dict(step=step, code=code, model=model.state_dict(), optimizer=training_manager.get_state()) |
| os.makedirs(f"logs/{run_id}", exist_ok=True) |
| torch.save(log, f"logs/{run_id}/state_step{step:06d}.pt") |
| if last_step: |
| break |
|
|
| |
| for idx in range(grad_accum_steps): |
| inputs, targets, cum_seqlens, bigram_inputs, bigram_cpu = train_loader.send(training_manager.train_loader_send_args) |
| training_manager.sparse_index_update(step, bigram_cpu) |
| loss = model(inputs, targets, cum_seqlens, bigram_inputs, training_manager.get_forward_args()).sum() * grad_scale |
| training_manager.sparse_index_share(step) |
| loss.backward() |
| del loss |
| training_manager.step_optimizers(step) |
|
|
| |
| approx_training_time_ms = training_time_ms + 1000 * (time.perf_counter() - t0) |
| print0(f"step:{step+1}/{train_steps} train_time:{approx_training_time_ms:.0f}ms step_avg:{approx_training_time_ms/(step + 1):.2f}ms", console=True) |
|
|
| if args.run_evals: |
| model.eval() |
| from evals import hellaswag |
| hellaswag.evaluate(model=model, |
| schedule_cfg=training_manager.get_forward_args(), |
| seq_len=args.val_batch_size // (grad_accum_steps * world_size), |
| get_bigram_hash=get_bigram_hash, |
| print0=print0) |
|
|
| print0(f"peak memory allocated: {torch.cuda.max_memory_allocated() // 1024 // 1024} MiB " |
| f"reserved: {torch.cuda.max_memory_reserved() // 1024 // 1024} MiB", console=True) |
| dist.destroy_process_group() |
|
|
|
|
| ---------------------------------------- |
| |
| ---------------------------------------- |
|
|
| import torch |
| import triton |
| import triton.language as tl |
| from triton.tools.tensor_descriptor import TensorDescriptor |
|
|
| |
| |
|
|
| @triton.jit |
| def _pid_to_block( |
| pid, |
| M, |
| BLOCK_SIZE_M: tl.constexpr, |
| BLOCK_SIZE_N: tl.constexpr, |
| GROUP_SIZE_M: tl.constexpr, |
| ): |
| |
| num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) |
| num_pid_n = tl.cdiv(M, BLOCK_SIZE_N) |
|
|
| |
| batch_idx = pid // (num_pid_m * num_pid_n) |
| pid = pid % (num_pid_m * num_pid_n) |
|
|
| |
| pid_m = pid // num_pid_n |
| pid_n = pid % num_pid_n |
| pid_m, pid_n = tl.swizzle2d(pid_m, pid_n, num_pid_m, num_pid_n, GROUP_SIZE_M) |
|
|
| m_idx = pid_m * BLOCK_SIZE_M |
| n_idx = pid_n * BLOCK_SIZE_N |
| return batch_idx, m_idx, n_idx |
|
|
| @triton.jit |
| def XXT_kernel( |
| A_ptr, C_ptr, |
| M, K, |
| a_stride_b, a_stride_r, a_stride_c, |
| c_stride_b, c_stride_r, c_stride_c, |
| BLOCK_SIZE_M: tl.constexpr, |
| BLOCK_SIZE_N: tl.constexpr, |
| BLOCK_SIZE_K: tl.constexpr, |
| GROUP_SIZE_M: tl.constexpr, |
| LOWER_UPPER: tl.constexpr, |
| ): |
| pid = tl.program_id(axis=0) |
| batch_idx, m_idx, n_idx = _pid_to_block( |
| pid, M, BLOCK_SIZE_M, BLOCK_SIZE_N, GROUP_SIZE_M |
| ) |
|
|
| |
| skip_block_below_diag = (LOWER_UPPER == 0) and (n_idx + BLOCK_SIZE_N <= m_idx) |
| skip_block_above_diag = (LOWER_UPPER != 0) and (m_idx + BLOCK_SIZE_M <= n_idx) |
| if skip_block_below_diag or skip_block_above_diag: |
| return |
|
|
| |
| A_ptr += batch_idx * a_stride_b |
| C_ptr += batch_idx * c_stride_b |
|
|
| |
| offs_m = (m_idx + tl.arange(0, BLOCK_SIZE_M)) % M |
| offs_n = (n_idx + tl.arange(0, BLOCK_SIZE_N)) % M |
| offs_k = tl.arange(0, BLOCK_SIZE_K) |
| |
| |
| |
| a_ptrs = A_ptr + (offs_m[:, None] * a_stride_r + offs_k[None, :] * a_stride_c) |
| |
| |
| at_ptrs = A_ptr + (offs_n[:, None] * a_stride_r + offs_k[None, :] * a_stride_c) |
|
|
| accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) |
|
|
| |
| for k in tl.range(0, tl.cdiv(K, BLOCK_SIZE_K)): |
| k_remaining = K - k * BLOCK_SIZE_K |
| a = tl.load(a_ptrs, mask=offs_k[None, :] < k_remaining, other=0.0) |
| at_temp = tl.load(at_ptrs, mask=offs_k[None, :] < k_remaining, other=0.0) |
| at = tl.trans(at_temp) |
| accumulator = tl.dot(a, at, accumulator) |
| a_ptrs += BLOCK_SIZE_K * a_stride_c |
| at_ptrs += BLOCK_SIZE_K * a_stride_c |
|
|
| out_dtype = C_ptr.dtype.element_ty |
| output = accumulator.to(out_dtype) |
|
|
| |
| offs_cm = m_idx + tl.arange(0, BLOCK_SIZE_M) |
| offs_cn = n_idx + tl.arange(0, BLOCK_SIZE_N) |
| c_ptrs = C_ptr + (offs_cm[:, None] * c_stride_r + offs_cn[None, :] * c_stride_c) |
| c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < M) |
| tl.store(c_ptrs, output, mask=c_mask) |
|
|
| |
| c_ptrs_t = C_ptr + (offs_cn[:, None] * c_stride_r + offs_cm[None, :] * c_stride_c) |
| c_mask_t = (offs_cn[:, None] < M) & (offs_cm[None, :] < M) |
| tl.store(c_ptrs_t, output.T, mask=c_mask_t) |
|
|
| def XXT(A: torch.Tensor, out: torch.Tensor): |
| """ |
| Launch Triton kernel to compute C = A @ A.T |
| """ |
| assert A.ndim == 2 or A.ndim == 3 |
| M, K = A.shape[-2:] |
| assert out.size(-2) == M, "Output matrix has incorrect shape" |
| assert out.size(-1) == M, "Output matrix has incorrect shape" |
|
|
| batch_size = A.size(0) if A.ndim == 3 else 1 |
| input_batch_stride = A.stride(0) if A.ndim == 3 else 0 |
| output_batch_stride = out.stride(0) if out.ndim == 3 else 0 |
|
|
| |
| if K == 768: |
| BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 128, 128, 64 |
| num_stages, num_warps = 4, 8 |
| else: |
| BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 64, 128, 128 |
| num_stages, num_warps = 4, 8 |
|
|
| grid = (batch_size * triton.cdiv(M, BLOCK_SIZE_M) * triton.cdiv(M, BLOCK_SIZE_N),) |
| XXT_kernel[grid]( |
| A_ptr=A, |
| C_ptr=out, |
| M=M, |
| K=K, |
| a_stride_b=input_batch_stride, |
| a_stride_r=A.stride(-2), |
| a_stride_c=A.stride(-1), |
| c_stride_b=output_batch_stride, |
| c_stride_r=out.stride(-2), |
| c_stride_c=out.stride(-1), |
| BLOCK_SIZE_M=BLOCK_SIZE_M, |
| BLOCK_SIZE_N=BLOCK_SIZE_N, |
| BLOCK_SIZE_K=BLOCK_SIZE_K, |
| GROUP_SIZE_M=8, |
| LOWER_UPPER=1, |
| num_stages=num_stages, |
| num_warps=num_warps, |
| ) |
| return out |
|
|
| |
| |
| |
|
|
| @triton.jit |
| def XTX_kernel( |
| A_ptr, C_ptr, |
| M, K, |
| a_stride_b, a_stride_r, a_stride_c, |
| c_stride_b, c_stride_r, c_stride_c, |
| BLOCK_SIZE_M: tl.constexpr, |
| BLOCK_SIZE_N: tl.constexpr, |
| BLOCK_SIZE_K: tl.constexpr, |
| GROUP_SIZE_M: tl.constexpr, |
| LOWER_UPPER: tl.constexpr, |
| ): |
| """ |
| Compute C = A.T @ A where A is (M, K) and C is (K, K). |
| This is the transpose variant of XXT for tall matrices. |
| |
| The output matrix C is symmetric, so we compute upper triangle and mirror. |
| We iterate over blocks of M (the reduction dimension after transpose). |
| """ |
| pid = tl.program_id(axis=0) |
| |
| batch_idx, k_idx, n_idx = _pid_to_block( |
| pid, K, BLOCK_SIZE_M, BLOCK_SIZE_N, GROUP_SIZE_M |
| ) |
|
|
| |
| skip_block_below_diag = (LOWER_UPPER == 0) and (n_idx + BLOCK_SIZE_N <= k_idx) |
| skip_block_above_diag = (LOWER_UPPER != 0) and (k_idx + BLOCK_SIZE_M <= n_idx) |
| if skip_block_below_diag or skip_block_above_diag: |
| return |
|
|
| |
| A_ptr += batch_idx * a_stride_b |
| C_ptr += batch_idx * c_stride_b |
|
|
| |
| |
| |
| |
| offs_k = (k_idx + tl.arange(0, BLOCK_SIZE_M)) % K |
| offs_n = (n_idx + tl.arange(0, BLOCK_SIZE_N)) % K |
| offs_m = tl.arange(0, BLOCK_SIZE_K) |
|
|
| |
| |
| at_ptrs = A_ptr + (offs_m[:, None] * a_stride_r + offs_k[None, :] * a_stride_c) |
| |
| a_ptrs = A_ptr + (offs_m[:, None] * a_stride_r + offs_n[None, :] * a_stride_c) |
|
|
| accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) |
|
|
| |
| for m in tl.range(0, tl.cdiv(M, BLOCK_SIZE_K)): |
| m_remaining = M - m * BLOCK_SIZE_K |
| |
| at = tl.load(at_ptrs, mask=offs_m[:, None] < m_remaining, other=0.0) |
| |
| a = tl.load(a_ptrs, mask=offs_m[:, None] < m_remaining, other=0.0) |
| |
| |
| accumulator = tl.dot(at.T, a, accumulator) |
| at_ptrs += BLOCK_SIZE_K * a_stride_r |
| a_ptrs += BLOCK_SIZE_K * a_stride_r |
|
|
| out_dtype = C_ptr.dtype.element_ty |
| output = accumulator.to(out_dtype) |
|
|
| |
| offs_ck = k_idx + tl.arange(0, BLOCK_SIZE_M) |
| offs_cn = n_idx + tl.arange(0, BLOCK_SIZE_N) |
| c_ptrs = C_ptr + (offs_ck[:, None] * c_stride_r + offs_cn[None, :] * c_stride_c) |
| c_mask = (offs_ck[:, None] < K) & (offs_cn[None, :] < K) |
| tl.store(c_ptrs, output, mask=c_mask) |
|
|
| |
| c_ptrs_t = C_ptr + (offs_cn[:, None] * c_stride_r + offs_ck[None, :] * c_stride_c) |
| c_mask_t = (offs_cn[:, None] < K) & (offs_ck[None, :] < K) |
| tl.store(c_ptrs_t, output.T, mask=c_mask_t) |
|
|
|
|
| def XTX(A: torch.Tensor, out: torch.Tensor): |
| """ |
| Launch Triton kernel to compute C = A.T @ A |
| |
| For tall matrices (M > K), this is more efficient than transposing |
| and using XXT because the intermediate products are smaller (K x K vs M x M). |
| |
| Args: |
| A: Input tensor of shape (M, K) or (batch, M, K) |
| out: Output tensor of shape (K, K) or (batch, K, K) |
| |
| Returns: |
| out: The same output tensor, filled with A.T @ A |
| """ |
| assert A.ndim == 2 or A.ndim == 3 |
| M, K = A.shape[-2:] |
| assert out.size(-2) == K, f"Output matrix has incorrect shape: expected ({K}, {K}), got {tuple(out.shape[-2:])}" |
| assert out.size(-1) == K, f"Output matrix has incorrect shape: expected ({K}, {K}), got {tuple(out.shape[-2:])}" |
|
|
| batch_size = A.size(0) if A.ndim == 3 else 1 |
| input_batch_stride = A.stride(0) if A.ndim == 3 else 0 |
| output_batch_stride = out.stride(0) if out.ndim == 3 else 0 |
|
|
| |
| if K == 768: |
| BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 128, 128, 64 |
| num_stages, num_warps = 4, 8 |
| else: |
| BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 64, 128, 128 |
| num_stages, num_warps = 4, 8 |
|
|
| grid = (batch_size * triton.cdiv(K, BLOCK_SIZE_M) * triton.cdiv(K, BLOCK_SIZE_N),) |
| XTX_kernel[grid]( |
| A_ptr=A, |
| C_ptr=out, |
| M=M, |
| K=K, |
| a_stride_b=input_batch_stride, |
| a_stride_r=A.stride(-2), |
| a_stride_c=A.stride(-1), |
| c_stride_b=output_batch_stride, |
| c_stride_r=out.stride(-2), |
| c_stride_c=out.stride(-1), |
| BLOCK_SIZE_M=BLOCK_SIZE_M, |
| BLOCK_SIZE_N=BLOCK_SIZE_N, |
| BLOCK_SIZE_K=BLOCK_SIZE_K, |
| GROUP_SIZE_M=8, |
| LOWER_UPPER=1, |
| num_stages=num_stages, |
| num_warps=num_warps, |
| ) |
| return out |
|
|
|
|
| @triton.jit |
| def ba_plus_cAA_kernel( |
| A_ptr, C_ptr, |
| M, |
| a_stride_b, a_stride_r, a_stride_c, |
| c_stride_b, c_stride_r, c_stride_c, |
| alpha, beta, |
| BLOCK_SIZE_M: tl.constexpr, |
| BLOCK_SIZE_N: tl.constexpr, |
| BLOCK_SIZE_K: tl.constexpr, |
| GROUP_SIZE_M: tl.constexpr, |
| LOWER_UPPER: tl.constexpr, |
| ): |
| |
| |
| pid = tl.program_id(axis=0) |
| batch_idx, m_idx, n_idx = _pid_to_block( |
| pid, M, BLOCK_SIZE_M, BLOCK_SIZE_N, GROUP_SIZE_M |
| ) |
|
|
| |
| skip_block_below_diag = (LOWER_UPPER == 0) and (n_idx + BLOCK_SIZE_N <= m_idx) |
| skip_block_above_diag = (LOWER_UPPER != 0) and (m_idx + BLOCK_SIZE_M <= n_idx) |
| if skip_block_below_diag or skip_block_above_diag: |
| return |
|
|
| |
| A_ptr += batch_idx * a_stride_b |
| C_ptr += batch_idx * c_stride_b |
|
|
| |
| offs_m = (m_idx + tl.arange(0, BLOCK_SIZE_M)) % M |
| offs_n = (n_idx + tl.arange(0, BLOCK_SIZE_N)) % M |
| offs_k = tl.arange(0, BLOCK_SIZE_K) |
| |
| |
| a_ptrs = A_ptr + (offs_m[:, None] * a_stride_r + offs_k[None, :] * a_stride_c) |
| at_ptrs = A_ptr + (offs_n[:, None] * a_stride_r + offs_k[None, :] * a_stride_c) |
|
|
| accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) |
|
|
| |
| for k in tl.range(0, tl.cdiv(M, BLOCK_SIZE_K)): |
| k_remaining = M - k * BLOCK_SIZE_K |
| a = tl.load(a_ptrs, mask=offs_k[None, :] < k_remaining, other=0.0) |
| at_temp = tl.load(at_ptrs, mask=offs_k[None, :] < k_remaining, other=0.0) |
| at = tl.trans(at_temp) |
| accumulator = tl.dot(a, at, accumulator) |
| a_ptrs += BLOCK_SIZE_K * a_stride_c |
| at_ptrs += BLOCK_SIZE_K * a_stride_c |
|
|
| |
| offs_am = m_idx + tl.arange(0, BLOCK_SIZE_M) |
| offs_an = n_idx + tl.arange(0, BLOCK_SIZE_N) |
| a_add_ptrs = A_ptr + (offs_am[:, None] * a_stride_r + offs_an[None, :] * a_stride_c) |
| a_add_mask = (offs_am[:, None] < M) & (offs_an[None, :] < M) |
| a_add = tl.load(a_add_ptrs, mask=a_add_mask, other=0.0).to(tl.float32) |
|
|
| |
| accumulator *= alpha |
| accumulator += a_add * beta |
|
|
| out_dtype = C_ptr.dtype.element_ty |
| output = accumulator.to(out_dtype) |
|
|
| |
| offs_cm = m_idx + tl.arange(0, BLOCK_SIZE_M) |
| offs_cn = n_idx + tl.arange(0, BLOCK_SIZE_N) |
| c_ptrs = C_ptr + (offs_cm[:, None] * c_stride_r + offs_cn[None, :] * c_stride_c) |
| c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < M) |
| tl.store(c_ptrs, output, mask=c_mask) |
|
|
| |
| c_ptrs_t = C_ptr + (offs_cn[:, None] * c_stride_r + offs_cm[None, :] * c_stride_c) |
| c_mask_t = (offs_cn[:, None] < M) & (offs_cm[None, :] < M) |
| tl.store(c_ptrs_t, output.T, mask=c_mask_t) |
|
|
| def ba_plus_cAA(A: torch.Tensor, alpha: float, beta: float, out: torch.Tensor): |
| """ |
| Launch Triton kernel to compute C = alpha * A @ A.T + beta * A |
| """ |
| assert A.ndim == 2 or A.ndim == 3 |
| M, K = A.shape[-2:] |
| assert M == K, "Input matrix must be square" |
| assert out.size(-2) == M |
| assert out.size(-1) == M |
|
|
| batch_size = A.size(0) if A.ndim == 3 else 1 |
| input_batch_stride = A.stride(0) if A.ndim == 3 else 0 |
| output_batch_stride = out.stride(0) if out.ndim == 3 else 0 |
|
|
| |
| BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 128, 128, 64 |
| num_stages, num_warps = 4, 8 |
|
|
| grid = (batch_size * triton.cdiv(M, BLOCK_SIZE_M) * triton.cdiv(M, BLOCK_SIZE_N),) |
| ba_plus_cAA_kernel[grid]( |
| A_ptr=A, |
| C_ptr=out, |
| M=M, |
| a_stride_b=input_batch_stride, |
| a_stride_r=A.stride(-2), |
| a_stride_c=A.stride(-1), |
| c_stride_b=output_batch_stride, |
| c_stride_r=out.stride(-2), |
| c_stride_c=out.stride(-1), |
| alpha=alpha, |
| beta=beta, |
| BLOCK_SIZE_M=BLOCK_SIZE_M, |
| BLOCK_SIZE_N=BLOCK_SIZE_N, |
| BLOCK_SIZE_K=BLOCK_SIZE_K, |
| GROUP_SIZE_M=8, |
| LOWER_UPPER=1, |
| num_stages=num_stages, |
| num_warps=num_warps, |
| ) |
| return out |
|
|
| |
| |
|
|
| @triton.jit |
| def linear_relu_square_kernel(a_desc, b_desc, c_desc, aux_desc, |
| M, N, K, |
| BLOCK_SIZE_M: tl.constexpr, |
| BLOCK_SIZE_N: tl.constexpr, |
| BLOCK_SIZE_K: tl.constexpr, |
| GROUP_SIZE_M: tl.constexpr, |
| NUM_SMS: tl.constexpr, |
| FORWARD: tl.constexpr, |
| ): |
| dtype = tl.bfloat16 |
| start_pid = tl.program_id(axis=0) |
| num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) |
| num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) |
| k_tiles = tl.cdiv(K, BLOCK_SIZE_K) |
| num_tiles = num_pid_m * num_pid_n |
|
|
| tile_id_c = start_pid - NUM_SMS |
| num_pid_in_group = GROUP_SIZE_M * num_pid_n |
|
|
| for tile_id in tl.range(start_pid, num_tiles, NUM_SMS, flatten=True): |
| pid_m = tile_id // num_pid_n |
| pid_n = tile_id % num_pid_n |
| offs_am = pid_m * BLOCK_SIZE_M |
| offs_bn = pid_n * BLOCK_SIZE_N |
|
|
| accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) |
| for ki in range(k_tiles): |
| offs_k = ki * BLOCK_SIZE_K |
| a = a_desc.load([offs_am, offs_k]) |
| b = b_desc.load([offs_bn, offs_k]) |
| accumulator = tl.dot(a, b.T, accumulator) |
|
|
| tile_id_c += NUM_SMS |
| pid_m = tile_id // num_pid_n |
| pid_n = tile_id % num_pid_n |
| offs_am_c = pid_m * BLOCK_SIZE_M |
| offs_bn_c = pid_n * BLOCK_SIZE_N |
|
|
| acc = tl.reshape(accumulator, (BLOCK_SIZE_M, 2, BLOCK_SIZE_N // 2)) |
| acc = tl.permute(acc, (0, 2, 1)) |
| acc0, acc1 = tl.split(acc) |
|
|
| c0 = acc0.to(dtype) |
| if not FORWARD: |
| c0_pre = aux_desc.load([offs_am_c, offs_bn_c]) |
| c0 = 2 * c0 * tl.where(c0_pre > 0, c0_pre, 0) |
|
|
| c_desc.store([offs_am_c, offs_bn_c], c0) |
|
|
| if FORWARD: |
| c0_post = tl.maximum(c0, 0) |
| c0_post = c0_post * c0_post |
| aux_desc.store([offs_am_c, offs_bn_c], c0_post) |
|
|
| c1 = acc1.to(dtype) |
| if not FORWARD: |
| c1_pre = aux_desc.load([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2]) |
| c1 = 2 * c1 * tl.where(c1_pre > 0, c1_pre, 0) |
|
|
| c_desc.store([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2], c1) |
|
|
| if FORWARD: |
| c1_post = tl.maximum(c1, 0) |
| c1_post = c1_post * c1_post |
| aux_desc.store([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2], c1_post) |
|
|
|
|
| def linear_relu_square(a, b, aux=None): |
| M, K = a.shape |
| N, K = b.shape |
| dtype = a.dtype |
|
|
| c = torch.empty((M, N), device=a.device, dtype=dtype) |
|
|
| FORWARD = False |
| if aux is None: |
| FORWARD = True |
| aux = torch.empty((M, N), device=a.device, dtype=dtype) |
|
|
| NUM_SMS = torch.cuda.get_device_properties("cuda").multi_processor_count |
|
|
| BLOCK_SIZE_M = 128 |
| BLOCK_SIZE_N = 256 |
| BLOCK_SIZE_K = 64 |
| num_stages = 4 if FORWARD else 3 |
| num_warps = 8 |
|
|
| a_desc = TensorDescriptor.from_tensor(a, [BLOCK_SIZE_M, BLOCK_SIZE_K]) |
| b_desc = TensorDescriptor.from_tensor(b, [BLOCK_SIZE_N, BLOCK_SIZE_K]) |
| c_desc = TensorDescriptor.from_tensor(c, [BLOCK_SIZE_M, BLOCK_SIZE_N // 2]) |
| aux_desc = TensorDescriptor.from_tensor(aux, [BLOCK_SIZE_M, BLOCK_SIZE_N // 2]) |
|
|
| def grid(META): |
| return (min( |
| NUM_SMS, |
| triton.cdiv(M, BLOCK_SIZE_M) * triton.cdiv(N, BLOCK_SIZE_N), |
| ), ) |
|
|
| linear_relu_square_kernel[grid]( |
| a_desc, b_desc, c_desc, aux_desc, |
| M, N, K, |
| BLOCK_SIZE_M=BLOCK_SIZE_M, |
| BLOCK_SIZE_N=BLOCK_SIZE_N, |
| BLOCK_SIZE_K=BLOCK_SIZE_K, |
| GROUP_SIZE_M=1, |
| NUM_SMS=NUM_SMS, |
| FORWARD=FORWARD, |
| num_stages=num_stages, |
| num_warps=num_warps |
| ) |
|
|
| if FORWARD: |
| return c, aux |
| else: |
| return c |
|
|
| class FusedLinearReLUSquareFunction(torch.autograd.Function): |
| @staticmethod |
| def forward(ctx, x, W1, W2): |
| pre, post = linear_relu_square(x.view((-1, x.shape[-1])), W1) |
| x3 = post @ W2 |
| ctx.save_for_backward(x, W1, W2, pre, post) |
| return x3.view(x.shape) |
|
|
| @staticmethod |
| def backward(ctx, grad_output): |
| x, W1, W2, pre, post = ctx.saved_tensors |
| dW2 = post.T @ grad_output |
| dpre = linear_relu_square(grad_output.view((-1, grad_output.shape[-1])), W2, aux=pre) |
| dW1 = dpre.T @ x |
| dx = dpre @ W1 |
| return dx.view(x.shape), dW1, dW2 |
|
|
|
|
| |
| |
| |
| |
| |
|
|
| @triton.jit |
| def _transpose_copy_kernel( |
| src_ptr, dst_ptr, |
| M, N, |
| src_stride_m, src_stride_n, |
| dst_stride_0, dst_stride_1, |
| BLOCK_M: tl.constexpr, |
| BLOCK_N: tl.constexpr, |
| ): |
| pid_m = tl.program_id(0) |
| pid_n = tl.program_id(1) |
|
|
| offs_m = (pid_m * BLOCK_M + tl.arange(0, BLOCK_M)).to(tl.int64) |
| offs_n = (pid_n * BLOCK_N + tl.arange(0, BLOCK_N)).to(tl.int64) |
|
|
| mask = (offs_m[:, None] < M) & (offs_n[None, :] < N) |
|
|
| |
| tile = tl.load( |
| src_ptr + offs_m[:, None] * src_stride_m + offs_n[None, :] * src_stride_n, |
| mask=mask, other=0.0, |
| ) |
|
|
| |
| mask_T = (offs_n[:, None] < N) & (offs_m[None, :] < M) |
| tl.store( |
| dst_ptr + offs_n[:, None] * dst_stride_0 + offs_m[None, :] * dst_stride_1, |
| tl.trans(tile), mask=mask_T, |
| ) |
|
|
|
|
| def transpose_copy(src: torch.Tensor, dst: torch.Tensor): |
| """Tiled transpose copy: dst = src.T where src is (M, N) and dst is (N, M). |
| |
| Uses a 64x128 tiled Triton kernel with coalesced reads AND writes, |
| achieving near memory-bandwidth-limited performance. |
| """ |
| assert src.ndim == 2 and dst.ndim == 2 |
| M, N = src.shape |
| assert dst.shape == (N, M), f"Expected dst shape ({N}, {M}), got {dst.shape}" |
|
|
| BLOCK_M, BLOCK_N = 64, 128 |
| grid = (triton.cdiv(M, BLOCK_M), triton.cdiv(N, BLOCK_N)) |
|
|
| _transpose_copy_kernel[grid]( |
| src, dst, |
| M, N, |
| src.stride(0), src.stride(1), |
| dst.stride(0), dst.stride(1), |
| BLOCK_M=BLOCK_M, |
| BLOCK_N=BLOCK_N, |
| num_warps=8, |
| num_stages=2, |
| ) |
|
|
|
|
| |
| |
| |
| |
| |
|
|
| @triton.jit |
| def _transpose_add_kernel( |
| src_ptr, dst_ptr, |
| M, N, |
| src_stride_m, src_stride_n, |
| dst_stride_0, dst_stride_1, |
| BLOCK_M: tl.constexpr, |
| BLOCK_N: tl.constexpr, |
| ): |
| pid_m = tl.program_id(0) |
| pid_n = tl.program_id(1) |
|
|
| offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) |
| offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N) |
|
|
| mask = (offs_m[:, None] < M) & (offs_n[None, :] < N) |
|
|
| |
| src_tile = tl.load( |
| src_ptr + offs_m[:, None] * src_stride_m + offs_n[None, :] * src_stride_n, |
| mask=mask, other=0.0, |
| ) |
|
|
| |
| mask_T = (offs_n[:, None] < N) & (offs_m[None, :] < M) |
| dst_ptrs = dst_ptr + offs_n[:, None] * dst_stride_0 + offs_m[None, :] * dst_stride_1 |
| dst_tile = tl.load(dst_ptrs, mask=mask_T, other=0.0) |
| tl.store(dst_ptrs, dst_tile + tl.trans(src_tile), mask=mask_T) |
|
|
|
|
| def transpose_add(src: torch.Tensor, dst: torch.Tensor): |
| """Tiled transpose-add: dst += src.T where src is (M, N) and dst is (N, M). |
| |
| Uses a 32x32 tiled Triton kernel with coalesced access on both src and dst, |
| replacing PyTorch's .add_(src.T) which has non-coalesced reads from the |
| transposed operand. |
| """ |
| assert src.ndim == 2 and dst.ndim == 2 |
| M, N = src.shape |
| assert dst.shape == (N, M), f"Expected dst shape ({N}, {M}), got {dst.shape}" |
|
|
| BLOCK_M, BLOCK_N = 32, 32 |
| grid = (triton.cdiv(M, BLOCK_M), triton.cdiv(N, BLOCK_N)) |
|
|
| _transpose_add_kernel[grid]( |
| src, dst, |
| M, N, |
| src.stride(0), src.stride(1), |
| dst.stride(0), dst.stride(1), |
| BLOCK_M=BLOCK_M, |
| BLOCK_N=BLOCK_N, |
| num_warps=4, |
| num_stages=2, |
| ) |
|
|
|
|
| CE_KERNEL_BLOCK_SIZE = 256 |
| CE_KERNEL_VOCAB_SIZE = 50304 |
|
|
| CE_KERNEL_DECLS = f""" |
| constexpr int VOCAB_SIZE = {CE_KERNEL_VOCAB_SIZE}; |
| constexpr int BLOCK_SIZE = {CE_KERNEL_BLOCK_SIZE}; |
| """ |
|
|
| CE_KERNEL_SOURCE = """ |
| #include <cuda_bf16.h> |
| #include <math_constants.h> |
| |
| #define __nv_fp8_e5m2 char |
| #define uint16_t unsigned short |
| #define uint8_t unsigned char |
| #define int64_t long long |
| |
| __device__ __forceinline__ __nv_fp8_e5m2 f32_to_fp8_e5m2(float x) { |
| uint16_t packed; |
| asm volatile( |
| "cvt.rn.satfinite.e5m2x2.f32 %0, %1, %2;" |
| : "=h"(packed) |
| : "f"(x), "f"(0.0f) |
| ); |
| __nv_fp8_e5m2 result; |
| *reinterpret_cast<uint8_t*>(&result) = (packed & (0xFF << 8)) >> 8; |
| return result; |
| } |
| |
| struct __align__(16) __nv_bfloat168 { |
| __nv_bfloat16 data[8]; |
| __device__ __nv_bfloat16& operator[](int i) { return data[i]; } |
| __device__ const __nv_bfloat16& operator[](int i) const { return data[i]; } |
| }; |
| |
| struct __align__(8) __nv_fp8_e5m28 { |
| __nv_fp8_e5m2 data[8]; |
| __device__ __nv_fp8_e5m2& operator[](int i) { return data[i]; } |
| __device__ const __nv_fp8_e5m2& operator[](int i) const { return data[i]; } |
| }; |
| |
| template<typename T> __device__ constexpr T CEIL_DIV(T a, T b) { return (a + b - 1) / b; } |
| |
| //__device__ float sigmoid(float x) { |
| // return 1.0f / (1.0f + __expf(-x)); |
| //} |
| __device__ float sigmoid(float x) { |
| return 0.5f + __tanhf(x * 0.5f) * 0.5f; |
| } |
| |
| extern "C" |
| __launch_bounds__(BLOCK_SIZE, 2) |
| __global__ void ce_fwd_bwd_kernel( |
| const __nv_bfloat16* __restrict__ logits, |
| const int64_t* __restrict__ targets, |
| const float* __restrict__ mtp_weights, |
| float* __restrict__ losses, |
| __nv_fp8_e5m2* grad_input, |
| int batch_size, |
| int n_predict, |
| double A_param, |
| double B_param, |
| double C_param, |
| double grad_s_param, |
| double grad_scale_param) |
| { |
| constexpr int VEC_WIDTH = 8; |
| constexpr int NUM_FULL_LOADS = VOCAB_SIZE / (BLOCK_SIZE * VEC_WIDTH); |
| constexpr int NUM_LOADS = CEIL_DIV(VOCAB_SIZE, BLOCK_SIZE * VEC_WIDTH); |
| |
| float A = (float)A_param; |
| float B = (float)B_param; |
| float C = (float)C_param; |
| float grad_s = (float)grad_s_param; |
| float grad_scale = (float)grad_scale_param; |
| |
| extern __shared__ __nv_bfloat16 smem[]; |
| |
| static_assert(VEC_WIDTH == 8); |
| |
| const __nv_bfloat16 *block_logit_ptr = logits + VOCAB_SIZE * blockIdx.x; |
| |
| float inv_C = 1 / C; |
| float B_div_C = B * inv_C; |
| float thread_max = -CUDART_INF_F; |
| |
| #pragma unroll 25 |
| for (int i = 0; i < NUM_LOADS; i++) { |
| int idx = i * BLOCK_SIZE * VEC_WIDTH + threadIdx.x * VEC_WIDTH; |
| if (i < NUM_FULL_LOADS || idx < VOCAB_SIZE) { |
| __nv_bfloat168 result = *(__nv_bfloat168*)(&block_logit_ptr[idx]); |
| __nv_bfloat168 result_sigmoid; |
| #pragma unroll |
| for (int k = 0; k < VEC_WIDTH; k++) { |
| float tmp = __bfloat162float(result[k]); |
| tmp = sigmoid(tmp * inv_C + B_div_C); |
| result_sigmoid[k] = __float2bfloat16(tmp); |
| tmp = A * tmp; |
| thread_max = max(tmp, thread_max); |
| } |
| *(__nv_bfloat168*)(&smem[idx]) = result_sigmoid; |
| } |
| } |
| |
| constexpr int NUM_WARPS = BLOCK_SIZE / 32; |
| int warp_id = threadIdx.x / 32; |
| __shared__ float block_maxs[NUM_WARPS]; |
| __shared__ float block_sums[NUM_WARPS]; |
| |
| for (int offset = 16; offset > 0; offset >>= 1) |
| thread_max = fmaxf(thread_max, __shfl_down_sync(0xFFFFFFFF, thread_max, offset)); |
| |
| if (threadIdx.x % 32 == 0) { |
| block_maxs[warp_id] = thread_max; |
| } |
| |
| __syncthreads(); |
| |
| float block_max = -CUDART_INF_F; |
| for (int i = 0; i < NUM_WARPS; i++) { |
| block_max = fmaxf(block_max, block_maxs[i]); |
| } |
| |
| float thread_sum = 0.0f; |
| #pragma unroll 2 |
| for (int i = 0; i < NUM_LOADS; i++) { |
| int idx = i * BLOCK_SIZE * VEC_WIDTH + threadIdx.x * VEC_WIDTH; |
| __nv_bfloat168 l; |
| if (i < NUM_FULL_LOADS || idx < VOCAB_SIZE) { |
| l = *(__nv_bfloat168*)(&smem[idx]); |
| } |
| #pragma unroll |
| for (int k = 0; k < VEC_WIDTH; k++) { |
| float tmp = A * __bfloat162float(l[k]); |
| tmp = __expf(tmp - block_max); |
| if (i < NUM_FULL_LOADS || idx < VOCAB_SIZE) { |
| thread_sum += tmp; |
| } |
| } |
| } |
| |
| for (int offset = 16; offset > 0; offset >>= 1) |
| thread_sum += __shfl_down_sync(0xFFFFFFFF, thread_sum, offset); |
| |
| if (threadIdx.x % 32 == 0) { |
| block_sums[warp_id] = thread_sum; |
| } |
| |
| __syncthreads(); |
| |
| float block_sum = 0.0f; |
| for (int i = 0; i < NUM_WARPS; i++) { |
| block_sum += block_sums[i]; |
| } |
| |
| float lse = block_max + __logf(block_sum); |
| |
| if (threadIdx.x == 0) { |
| float total_loss = 0.0f; |
| for (int k = 0; k < n_predict; k++) { |
| int64_t target_idx = blockIdx.x + k; |
| if (target_idx < batch_size) { |
| float weight = mtp_weights[k]; |
| int64_t target = targets[target_idx]; |
| if (target >= 0 && target < VOCAB_SIZE) { |
| float z_target = A * __bfloat162float(smem[target]); |
| total_loss += weight * (lse - z_target); |
| } |
| } |
| } |
| losses[blockIdx.x] = total_loss; |
| } |
| |
| float S_w = 0.0f; |
| |
| for (int i = 0; i < n_predict; i++) { |
| S_w += mtp_weights[i]; |
| } |
| |
| #pragma unroll 4 |
| for (int i = 0; i < NUM_LOADS; i++) { |
| int idx = i * BLOCK_SIZE * VEC_WIDTH + threadIdx.x * VEC_WIDTH; |
| __nv_fp8_e5m28 result; |
| |
| if (i < NUM_FULL_LOADS || idx < VOCAB_SIZE) { |
| __nv_bfloat168 sigmoid_us = *(__nv_bfloat168*)(&smem[idx]); |
| #pragma unroll |
| for (int j = 0; j < VEC_WIDTH; j++) { |
| float sigmoid_u = __bfloat162float(sigmoid_us[j]); |
| float z = A * sigmoid_u; |
| float p = __expf(z - lse); |
| |
| float term1 = S_w * p; |
| float term2 = 0.0f; |
| |
| float grad_z = term1 - term2; |
| float grad_x = grad_scale * (1.0f / C * A) * (1.0f / grad_s) * grad_z * sigmoid_u * (1.0f - sigmoid_u); |
| auto result_tmp = f32_to_fp8_e5m2(grad_x); |
| result[j] = *reinterpret_cast<__nv_fp8_e5m2*>(&result_tmp); |
| } |
| *(__nv_fp8_e5m28*)(&grad_input[blockIdx.x * VOCAB_SIZE + idx]) = result; |
| } |
| } |
| |
| __syncthreads(); |
| |
| if (threadIdx.x < n_predict && blockIdx.x + threadIdx.x < batch_size) { |
| int i = threadIdx.x; |
| int64_t target = targets[blockIdx.x + i]; |
| |
| float sigmoid_u = __bfloat162float(smem[target]); |
| float z = A * sigmoid_u; |
| float p = __expf(z - lse); |
| |
| float term1 = S_w * p; |
| float term2 = 0.0f; |
| |
| #pragma unroll |
| for (int k = 0; k < 3; k++) { |
| int64_t target_idx = blockIdx.x + k; |
| if (target_idx < batch_size && k < n_predict) { |
| if (targets[target_idx] == target) { |
| term2 += mtp_weights[k]; |
| } |
| } |
| } |
| |
| float grad_z = term1 - term2; |
| float grad_x = grad_scale * (1.0f / C * A) * (1.0f / grad_s) * grad_z * sigmoid_u * (1.0f - sigmoid_u); |
| auto result_tmp = f32_to_fp8_e5m2(grad_x); |
| auto result = *reinterpret_cast<__nv_fp8_e5m2*>(&result_tmp); |
| grad_input[blockIdx.x * VOCAB_SIZE + target] = result; |
| } |
| } |
| """ |
|
|
| ce_fwd_bwd_kernel = torch.cuda._compile_kernel( |
| CE_KERNEL_DECLS + CE_KERNEL_SOURCE, |
| "ce_fwd_bwd_kernel", |
| compute_capability="90", |
| cuda_include_dirs=['/home/ubuntu/modded-nanogpt/.venv/lib/python3.14/site-packages/triton/backends/nvidia/include', '/home/ubuntu/modded-nanogpt/.venv/lib/python3.14/site-packages/nvidia/cuda_runtime/include'], |
| nvcc_options=["-lineinfo", "--use_fast_math"], |
| ) |
| ce_fwd_bwd_kernel.set_shared_memory_config(CE_KERNEL_VOCAB_SIZE * 2) |
|
|
| @torch.library.custom_op("nanogpt::ce_fwd_bwd", mutates_args={"losses", "grad_input"}) |
| def ce_fwd_bwd( |
| logits: torch.Tensor, |
| targets: torch.Tensor, |
| mtp_weights: torch.Tensor, |
| losses: torch.Tensor, |
| grad_input: torch.Tensor, |
| n_rows: int, |
| n_predict: int, |
| A: float, |
| B: float, |
| C: float, |
| grad_s: float, |
| grad_scale: float, |
| ) -> None: |
| grid = (n_rows, 1, 1) |
| ce_fwd_bwd_kernel( |
| grid, |
| (CE_KERNEL_BLOCK_SIZE, 1, 1), |
| (logits, targets, mtp_weights, losses, grad_input, |
| n_rows, n_predict, A, B, C, grad_s, grad_scale), |
| shared_mem=CE_KERNEL_VOCAB_SIZE * 2, |
| ) |
|
|
| class FusedSoftcappedCrossEntropy(torch.autograd.Function): |
| @staticmethod |
| def forward(ctx, x, targets, mtp_weights, lm_head_weight, x_s, w_s, grad_s, grad_scale, A=23.0, B=5.0, C=7.5): |
|
|
| x_f8 = x.div(x_s).to(torch.float8_e4m3fn) |
| w_f8 = lm_head_weight.div(w_s).to(torch.float8_e4m3fn) |
|
|
| w_f8_col_major = w_f8.T.contiguous().T |
|
|
| logits = torch._scaled_mm( |
| x_f8, |
| w_f8_col_major, |
| out_dtype=torch.bfloat16, |
| scale_a=x.new_tensor(x_s, dtype=torch.float32), |
| scale_b=x.new_tensor(w_s, dtype=torch.float32), |
| use_fast_accum=True, |
| ) |
|
|
| n_rows, n_cols = logits.shape |
| if mtp_weights is None: |
| mtp_weights = torch.tensor([1.0], device=logits.device, dtype=torch.float32) |
| n_predict = mtp_weights.shape[0] |
|
|
| losses = torch.empty(n_rows, dtype=torch.float32, device=logits.device) |
| lse = torch.empty(n_rows, dtype=torch.float32, device=logits.device) |
|
|
| logits = logits.contiguous() |
| targets = targets.contiguous() |
| mtp_weights = mtp_weights.contiguous() |
|
|
| grad_input = torch.empty((n_rows, n_cols), dtype=torch.float8_e5m2, device=logits.device) |
|
|
| ce_fwd_bwd(logits, targets, mtp_weights, losses, grad_input, |
| n_rows, n_predict, A, B, C, grad_s, grad_scale) |
|
|
| ctx.save_for_backward(logits, targets, mtp_weights, lse, x, lm_head_weight, x_f8, w_f8, grad_input) |
| ctx.params = (A, B, C, x_s, w_s, grad_s) |
| return losses |
|
|
| @staticmethod |
| def backward(ctx, grad_output): |
| logits, targets, mtp_weights, lse, x, lm_head_weight, x_f8, w_f8, grad_input = ctx.saved_tensors |
| A, B, C, x_s, w_s, grad_s = ctx.params |
| n_rows, n_cols = logits.shape |
| n_predict = mtp_weights.shape[0] |
|
|
| grad_output = grad_output.contiguous() |
|
|
| x_scale = grad_input.new_tensor(x_s, dtype=torch.float32) |
| w_scale = grad_input.new_tensor(w_s, dtype=torch.float32) |
| grad_scale = grad_input.new_tensor(grad_s, dtype=torch.float32) |
|
|
| grad_x = torch._scaled_mm( |
| grad_input, |
| w_f8.T, |
| out_dtype=torch.bfloat16, |
| scale_a=grad_scale, |
| scale_b=w_scale, |
| use_fast_accum=False, |
| ) |
|
|
| x_f8_T = torch.empty((x_f8.shape[1], x_f8.shape[0]), dtype=x_f8.dtype, device=x_f8.device) |
| transpose_copy(x_f8, x_f8_T) |
|
|
| grad_input_T = torch.empty((n_cols, n_rows), dtype=grad_input.dtype, device=grad_input.device) |
| transpose_copy(grad_input, grad_input_T) |
|
|
| grad_w = torch._scaled_mm( |
| x_f8_T, |
| grad_input_T.T, |
| out_dtype=torch.float32, |
| scale_a=x_scale, |
| scale_b=grad_scale, |
| use_fast_accum=False, |
| ) |
|
|
| return grad_x, None, None, grad_w, None, None, None |
|
|
|
|
| ==================================================================================================== |
| Running Python 3.14.4 (main, Apr 8 2026, 04:02:31) [GCC 15.2.0] |
| Running PyTorch 2.10.0+cu128 compiled for CUDA 12.8 |
| Running Triton version 3.6.0 |
| Mon Jun 15 01:22:16 2026 |
| +-----------------------------------------------------------------------------------------+ |
| | NVIDIA-SMI 595.71.05 Driver Version: 595.71.05 CUDA Version: 13.2 | |
| +-----------------------------------------+------------------------+----------------------+ |
| | GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC | |
| | Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. | |
| | | | MIG M. | |
| |=========================================+========================+======================| |
| | 0 NVIDIA H100 80GB HBM3 Off | 00000000:33:00.0 Off | 0 | |
| | N/A 41C P0 125W / 700W | 1189MiB / 81559MiB | 8% Default | |
| | | | Disabled | |
| +-----------------------------------------+------------------------+----------------------+ |
|
|
| +-----------------------------------------------------------------------------------------+ |
| | Processes: | |
| | GPU GI CI PID Type Process name GPU Memory | |
| | ID ID Usage | |
| |=========================================================================================| |
| | 0 N/A N/A 188886 C ...ded-nanogpt/.venv/bin/python3 1180MiB | |
| +-----------------------------------------------------------------------------------------+ |
|
|
| ==================================================================================================== |
| Compiling model and warming up kernels (~7 minutes on first execution) |
| Sampling steps [0, 1, 4398, 4399, 4400, 4401, 8798, 8799, 8800, 8801, 13198, 13199, 13200, 13201] for warmup |
| Resetting Model |
| step:0/13210 val_loss:10.4072 train_time:0ms step_avg:0.03ms |
| step:1/13210 train_time:247ms step_avg:247.35ms |
| step:2/13210 train_time:475ms step_avg:237.64ms |
| step:3/13210 train_time:704ms step_avg:234.60ms |
| step:4/13210 train_time:932ms step_avg:232.91ms |
| step:5/13210 train_time:1168ms step_avg:233.62ms |
| step:6/13210 train_time:1399ms step_avg:233.19ms |
| step:7/13210 train_time:1630ms step_avg:232.86ms |
| step:8/13210 train_time:1859ms step_avg:232.42ms |
| step:9/13210 train_time:2093ms step_avg:232.52ms |
| step:10/13210 train_time:2327ms step_avg:232.71ms |
| step:11/13210 train_time:2558ms step_avg:232.58ms |
| step:12/13210 train_time:2789ms step_avg:232.40ms |
| step:13/13210 train_time:3020ms step_avg:232.29ms |
| step:14/13210 train_time:3252ms step_avg:232.31ms |
| step:15/13210 train_time:3484ms step_avg:232.26ms |
| step:16/13210 train_time:3943ms step_avg:246.42ms |
| step:17/13210 train_time:4171ms step_avg:245.37ms |
| step:18/13210 train_time:4401ms step_avg:244.50ms |
| step:19/13210 train_time:4633ms step_avg:243.85ms |
| step:20/13210 train_time:4872ms step_avg:243.61ms |
| step:21/13210 train_time:5103ms step_avg:243.00ms |
| step:22/13210 train_time:5334ms step_avg:242.48ms |
| step:23/13210 train_time:5564ms step_avg:241.93ms |
| step:24/13210 train_time:5796ms step_avg:241.52ms |
| step:25/13210 train_time:6029ms step_avg:241.15ms |
| step:26/13210 train_time:6261ms step_avg:240.81ms |
| step:27/13210 train_time:6491ms step_avg:240.42ms |
| step:28/13210 train_time:6722ms step_avg:240.08ms |
| step:29/13210 train_time:6956ms step_avg:239.85ms |
| step:30/13210 train_time:7187ms step_avg:239.57ms |
| step:31/13210 train_time:7419ms step_avg:239.32ms |
| step:32/13210 train_time:7649ms step_avg:239.03ms |
| step:33/13210 train_time:7880ms step_avg:238.79ms |
| step:34/13210 train_time:8111ms step_avg:238.56ms |
| step:35/13210 train_time:8340ms step_avg:238.29ms |
| step:36/13210 train_time:8571ms step_avg:238.08ms |
| step:37/13210 train_time:8801ms step_avg:237.87ms |
| step:38/13210 train_time:9033ms step_avg:237.71ms |
| step:39/13210 train_time:9264ms step_avg:237.55ms |
| step:40/13210 train_time:9495ms step_avg:237.38ms |
| step:41/13210 train_time:9726ms step_avg:237.22ms |
| step:42/13210 train_time:9955ms step_avg:237.03ms |
| step:43/13210 train_time:10187ms step_avg:236.90ms |
| step:44/13210 train_time:10415ms step_avg:236.71ms |
| step:45/13210 train_time:10644ms step_avg:236.54ms |
| step:46/13210 train_time:10875ms step_avg:236.41ms |
| step:47/13210 train_time:11107ms step_avg:236.32ms |
| step:48/13210 train_time:11338ms step_avg:236.21ms |
| step:49/13210 train_time:11572ms step_avg:236.16ms |
| step:50/13210 train_time:11799ms step_avg:235.99ms |
| step:51/13210 train_time:12030ms step_avg:235.89ms |
| step:52/13210 train_time:12260ms step_avg:235.77ms |
| step:53/13210 train_time:12493ms step_avg:235.71ms |
| step:54/13210 train_time:12723ms step_avg:235.61ms |
| step:55/13210 train_time:12952ms step_avg:235.50ms |
| step:56/13210 train_time:13181ms step_avg:235.38ms |
| step:57/13210 train_time:13412ms step_avg:235.30ms |
| step:58/13210 train_time:13643ms step_avg:235.22ms |
| step:59/13210 train_time:13996ms step_avg:237.22ms |
| step:60/13210 train_time:14311ms step_avg:238.52ms |
| step:61/13210 train_time:14539ms step_avg:238.34ms |
| step:62/13210 train_time:14765ms step_avg:238.14ms |
| step:63/13210 train_time:14992ms step_avg:237.96ms |
| step:64/13210 train_time:15228ms step_avg:237.94ms |
| step:65/13210 train_time:15458ms step_avg:237.81ms |
| step:66/13210 train_time:15688ms step_avg:237.69ms |
| step:67/13210 train_time:15918ms step_avg:237.58ms |
| step:68/13210 train_time:16148ms step_avg:237.48ms |
| step:69/13210 train_time:16382ms step_avg:237.41ms |
| step:70/13210 train_time:16611ms step_avg:237.30ms |
| step:71/13210 train_time:16839ms step_avg:237.16ms |
| step:72/13210 train_time:17070ms step_avg:237.08ms |
| step:73/13210 train_time:17302ms step_avg:237.01ms |
| step:74/13210 train_time:17533ms step_avg:236.93ms |
| step:75/13210 train_time:17763ms step_avg:236.83ms |
| step:76/13210 train_time:17990ms step_avg:236.70ms |
| step:77/13210 train_time:18220ms step_avg:236.62ms |
| step:78/13210 train_time:18448ms step_avg:236.51ms |
| step:79/13210 train_time:18678ms step_avg:236.43ms |
| step:80/13210 train_time:18906ms step_avg:236.33ms |
| step:81/13210 train_time:19137ms step_avg:236.26ms |
| step:82/13210 train_time:19367ms step_avg:236.18ms |
| step:83/13210 train_time:19597ms step_avg:236.10ms |
| step:84/13210 train_time:19825ms step_avg:236.01ms |
| step:85/13210 train_time:20054ms step_avg:235.93ms |
| step:86/13210 train_time:20287ms step_avg:235.89ms |
| step:87/13210 train_time:20517ms step_avg:235.83ms |
| step:88/13210 train_time:20746ms step_avg:235.75ms |
| step:89/13210 train_time:20976ms step_avg:235.69ms |
| step:90/13210 train_time:21204ms step_avg:235.60ms |
| step:91/13210 train_time:21436ms step_avg:235.56ms |
| step:92/13210 train_time:21665ms step_avg:235.49ms |
| step:93/13210 train_time:21895ms step_avg:235.43ms |
| step:94/13210 train_time:22124ms step_avg:235.36ms |
| step:95/13210 train_time:22356ms step_avg:235.33ms |
| step:96/13210 train_time:22588ms step_avg:235.29ms |
| step:97/13210 train_time:22817ms step_avg:235.23ms |
| step:98/13210 train_time:23048ms step_avg:235.18ms |
| step:99/13210 train_time:23280ms step_avg:235.15ms |
| step:100/13210 train_time:23509ms step_avg:235.09ms |
| step:101/13210 train_time:23738ms step_avg:235.03ms |
| step:102/13210 train_time:23969ms step_avg:234.99ms |
| step:103/13210 train_time:24200ms step_avg:234.95ms |
| step:104/13210 train_time:24430ms step_avg:234.91ms |
| step:105/13210 train_time:24659ms step_avg:234.85ms |
| step:106/13210 train_time:24887ms step_avg:234.78ms |
| step:107/13210 train_time:25117ms step_avg:234.73ms |
| step:108/13210 train_time:25344ms step_avg:234.67ms |
| step:109/13210 train_time:25575ms step_avg:234.63ms |
| step:110/13210 train_time:25804ms step_avg:234.58ms |
| step:111/13210 train_time:26036ms step_avg:234.55ms |
| step:112/13210 train_time:26267ms step_avg:234.52ms |
| step:113/13210 train_time:26496ms step_avg:234.48ms |
| step:114/13210 train_time:26728ms step_avg:234.45ms |
| step:115/13210 train_time:26957ms step_avg:234.41ms |
| step:116/13210 train_time:27188ms step_avg:234.38ms |
| step:117/13210 train_time:27418ms step_avg:234.34ms |
| step:118/13210 train_time:27649ms step_avg:234.31ms |
| step:119/13210 train_time:27878ms step_avg:234.27ms |
| step:120/13210 train_time:28107ms step_avg:234.23ms |
| step:121/13210 train_time:28335ms step_avg:234.18ms |
| step:122/13210 train_time:28566ms step_avg:234.14ms |
| step:123/13210 train_time:28795ms step_avg:234.10ms |
| step:124/13210 train_time:29026ms step_avg:234.08ms |
| step:125/13210 train_time:29255ms step_avg:234.04ms |
| step:126/13210 train_time:29487ms step_avg:234.02ms |
| step:127/13210 train_time:29717ms step_avg:233.99ms |
| step:128/13210 train_time:29947ms step_avg:233.96ms |
| step:129/13210 train_time:30179ms step_avg:233.95ms |
| step:130/13210 train_time:30411ms step_avg:233.93ms |
| step:131/13210 train_time:30640ms step_avg:233.89ms |
| step:132/13210 train_time:30869ms step_avg:233.85ms |
| step:133/13210 train_time:31099ms step_avg:233.82ms |
| step:134/13210 train_time:31329ms step_avg:233.80ms |
| step:135/13210 train_time:31558ms step_avg:233.76ms |
| step:136/13210 train_time:31788ms step_avg:233.74ms |
| step:137/13210 train_time:32018ms step_avg:233.70ms |
| step:138/13210 train_time:32247ms step_avg:233.67ms |
| step:139/13210 train_time:32478ms step_avg:233.66ms |
| step:140/13210 train_time:32708ms step_avg:233.63ms |
| step:141/13210 train_time:32937ms step_avg:233.60ms |
| step:142/13210 train_time:33167ms step_avg:233.57ms |
| step:143/13210 train_time:33399ms step_avg:233.56ms |
| step:144/13210 train_time:33630ms step_avg:233.54ms |
| step:145/13210 train_time:33859ms step_avg:233.51ms |
| step:146/13210 train_time:34088ms step_avg:233.48ms |
| step:147/13210 train_time:34317ms step_avg:233.45ms |
| step:148/13210 train_time:34546ms step_avg:233.42ms |
| step:149/13210 train_time:34776ms step_avg:233.39ms |
| step:150/13210 train_time:35004ms step_avg:233.36ms |
| step:151/13210 train_time:35233ms step_avg:233.33ms |
| step:152/13210 train_time:35464ms step_avg:233.32ms |
| step:153/13210 train_time:35694ms step_avg:233.30ms |
| step:154/13210 train_time:35924ms step_avg:233.27ms |
| step:155/13210 train_time:36155ms step_avg:233.26ms |
| step:156/13210 train_time:36386ms step_avg:233.25ms |
| step:157/13210 train_time:36616ms step_avg:233.22ms |
| step:158/13210 train_time:36842ms step_avg:233.18ms |
| step:159/13210 train_time:37072ms step_avg:233.16ms |
| step:160/13210 train_time:37300ms step_avg:233.12ms |
| step:161/13210 train_time:37831ms step_avg:234.97ms |
| step:162/13210 train_time:38058ms step_avg:234.92ms |
| step:163/13210 train_time:38285ms step_avg:234.88ms |
| step:164/13210 train_time:38511ms step_avg:234.82ms |
| step:165/13210 train_time:38746ms step_avg:234.82ms |
| step:166/13210 train_time:38976ms step_avg:234.80ms |
| step:167/13210 train_time:39205ms step_avg:234.76ms |
| step:168/13210 train_time:39432ms step_avg:234.71ms |
| step:169/13210 train_time:39663ms step_avg:234.69ms |
| step:170/13210 train_time:39894ms step_avg:234.67ms |
| step:171/13210 train_time:40122ms step_avg:234.63ms |
| step:172/13210 train_time:40348ms step_avg:234.58ms |
| step:173/13210 train_time:40578ms step_avg:234.55ms |
| step:174/13210 train_time:40808ms step_avg:234.53ms |
| step:175/13210 train_time:41038ms step_avg:234.50ms |
| step:176/13210 train_time:41266ms step_avg:234.47ms |
| step:177/13210 train_time:41494ms step_avg:234.43ms |
| step:178/13210 train_time:41727ms step_avg:234.42ms |
| step:179/13210 train_time:41956ms step_avg:234.39ms |
| step:180/13210 train_time:42183ms step_avg:234.35ms |
| step:181/13210 train_time:42411ms step_avg:234.32ms |
| step:182/13210 train_time:42642ms step_avg:234.30ms |
| step:183/13210 train_time:42872ms step_avg:234.27ms |
| step:184/13210 train_time:43100ms step_avg:234.24ms |
| step:185/13210 train_time:43328ms step_avg:234.21ms |
| step:186/13210 train_time:43556ms step_avg:234.17ms |
| step:187/13210 train_time:43788ms step_avg:234.16ms |
| step:188/13210 train_time:44018ms step_avg:234.14ms |
| step:189/13210 train_time:44247ms step_avg:234.11ms |
| step:190/13210 train_time:44475ms step_avg:234.08ms |
| step:191/13210 train_time:45485ms step_avg:238.14ms |
| step:192/13210 train_time:45710ms step_avg:238.07ms |
| step:193/13210 train_time:45936ms step_avg:238.01ms |
| step:194/13210 train_time:46162ms step_avg:237.95ms |
| step:195/13210 train_time:46395ms step_avg:237.92ms |
| step:196/13210 train_time:46623ms step_avg:237.87ms |
| step:197/13210 train_time:46850ms step_avg:237.82ms |
| step:198/13210 train_time:47099ms step_avg:237.87ms |
| step:199/13210 train_time:47361ms step_avg:237.99ms |
| step:200/13210 train_time:47631ms step_avg:238.16ms |
| step:201/13210 train_time:47875ms step_avg:238.18ms |
| step:202/13210 train_time:48111ms step_avg:238.18ms |
| step:203/13210 train_time:48346ms step_avg:238.16ms |
| step:204/13210 train_time:48582ms step_avg:238.15ms |
| step:205/13210 train_time:48818ms step_avg:238.14ms |
| step:206/13210 train_time:49057ms step_avg:238.14ms |
| step:207/13210 train_time:49290ms step_avg:238.12ms |
| step:208/13210 train_time:49526ms step_avg:238.11ms |
| step:209/13210 train_time:49763ms step_avg:238.10ms |
| step:210/13210 train_time:50003ms step_avg:238.11ms |
| step:211/13210 train_time:50239ms step_avg:238.10ms |
| step:212/13210 train_time:50474ms step_avg:238.08ms |
| step:213/13210 train_time:50711ms step_avg:238.08ms |
| step:214/13210 train_time:50951ms step_avg:238.09ms |
| step:215/13210 train_time:51205ms step_avg:238.16ms |
| step:216/13210 train_time:51456ms step_avg:238.22ms |
| step:217/13210 train_time:51710ms step_avg:238.30ms |
| step:218/13210 train_time:51962ms step_avg:238.36ms |
| step:219/13210 train_time:52215ms step_avg:238.43ms |
| step:220/13210 train_time:52474ms step_avg:238.52ms |
| step:221/13210 train_time:52733ms step_avg:238.61ms |
| step:222/13210 train_time:53000ms step_avg:238.74ms |
| step:223/13210 train_time:53273ms step_avg:238.89ms |
| step:224/13210 train_time:53543ms step_avg:239.03ms |
| step:225/13210 train_time:53817ms step_avg:239.18ms |
| step:226/13210 train_time:54086ms step_avg:239.32ms |
| step:227/13210 train_time:54351ms step_avg:239.43ms |
| step:228/13210 train_time:54621ms step_avg:239.57ms |
| step:229/13210 train_time:54890ms step_avg:239.69ms |
| step:230/13210 train_time:55127ms step_avg:239.68ms |
| step:231/13210 train_time:55365ms step_avg:239.67ms |
| step:232/13210 train_time:55603ms step_avg:239.67ms |
| step:233/13210 train_time:55844ms step_avg:239.67ms |
| step:234/13210 train_time:56084ms step_avg:239.68ms |
| step:235/13210 train_time:56330ms step_avg:239.70ms |
| step:236/13210 train_time:56578ms step_avg:239.74ms |
| step:237/13210 train_time:56826ms step_avg:239.77ms |
| step:238/13210 train_time:57071ms step_avg:239.79ms |
| step:239/13210 train_time:57318ms step_avg:239.82ms |
| step:240/13210 train_time:57580ms step_avg:239.92ms |
| step:241/13210 train_time:57817ms step_avg:239.90ms |
| step:242/13210 train_time:58054ms step_avg:239.89ms |
| step:243/13210 train_time:58293ms step_avg:239.89ms |
| step:244/13210 train_time:58535ms step_avg:239.90ms |
| step:245/13210 train_time:58774ms step_avg:239.90ms |
| step:246/13210 train_time:59011ms step_avg:239.88ms |
| step:247/13210 train_time:59239ms step_avg:239.83ms |
| step:248/13210 train_time:59467ms step_avg:239.79ms |
| step:249/13210 train_time:59697ms step_avg:239.75ms |
| step:250/13210 train_time:59927ms step_avg:239.71ms |
| step:250/13210 val_loss:10.2278 train_time:59963ms step_avg:239.85ms |
| step:251/13210 train_time:60156ms step_avg:239.67ms |
| step:252/13210 train_time:60384ms step_avg:239.62ms |
| step:253/13210 train_time:60611ms step_avg:239.57ms |
| step:254/13210 train_time:60840ms step_avg:239.53ms |
| step:255/13210 train_time:61072ms step_avg:239.50ms |
| step:256/13210 train_time:61301ms step_avg:239.46ms |
| step:257/13210 train_time:61531ms step_avg:239.42ms |
| step:258/13210 train_time:61759ms step_avg:239.37ms |
| step:259/13210 train_time:61991ms step_avg:239.35ms |
| step:260/13210 train_time:62219ms step_avg:239.31ms |
| step:261/13210 train_time:62448ms step_avg:239.26ms |
| step:262/13210 train_time:62674ms step_avg:239.21ms |
| step:263/13210 train_time:62904ms step_avg:239.18ms |
| step:264/13210 train_time:63133ms step_avg:239.14ms |
| step:265/13210 train_time:63363ms step_avg:239.10ms |
| step:266/13210 train_time:63591ms step_avg:239.06ms |
| step:267/13210 train_time:63819ms step_avg:239.02ms |
| step:268/13210 train_time:64048ms step_avg:238.99ms |
| step:269/13210 train_time:64277ms step_avg:238.95ms |
| step:270/13210 train_time:64506ms step_avg:238.91ms |
| step:271/13210 train_time:64735ms step_avg:238.87ms |
| step:272/13210 train_time:64963ms step_avg:238.84ms |
| step:273/13210 train_time:65194ms step_avg:238.80ms |
| step:274/13210 train_time:65421ms step_avg:238.76ms |
| step:275/13210 train_time:65649ms step_avg:238.73ms |
| step:276/13210 train_time:65879ms step_avg:238.69ms |
| step:277/13210 train_time:66109ms step_avg:238.66ms |
| step:278/13210 train_time:66337ms step_avg:238.62ms |
| step:279/13210 train_time:66566ms step_avg:238.59ms |
| step:280/13210 train_time:66795ms step_avg:238.56ms |
| step:281/13210 train_time:67025ms step_avg:238.52ms |
| step:282/13210 train_time:67256ms step_avg:238.50ms |
| step:283/13210 train_time:67486ms step_avg:238.46ms |
| step:284/13210 train_time:67713ms step_avg:238.43ms |
| step:285/13210 train_time:67941ms step_avg:238.39ms |
| step:286/13210 train_time:68170ms step_avg:238.36ms |
| step:287/13210 train_time:68398ms step_avg:238.32ms |
| step:288/13210 train_time:68627ms step_avg:238.29ms |
| step:289/13210 train_time:68857ms step_avg:238.26ms |
| step:290/13210 train_time:69087ms step_avg:238.23ms |
| step:291/13210 train_time:69316ms step_avg:238.20ms |
| step:292/13210 train_time:69544ms step_avg:238.17ms |
| step:293/13210 train_time:69773ms step_avg:238.13ms |
| step:294/13210 train_time:70001ms step_avg:238.10ms |
| step:295/13210 train_time:70230ms step_avg:238.07ms |
| step:296/13210 train_time:70461ms step_avg:238.04ms |
| step:297/13210 train_time:70689ms step_avg:238.01ms |
| step:298/13210 train_time:70917ms step_avg:237.98ms |
| step:299/13210 train_time:71147ms step_avg:237.95ms |
| step:300/13210 train_time:71375ms step_avg:237.92ms |
| step:301/13210 train_time:71605ms step_avg:237.89ms |
| step:302/13210 train_time:71834ms step_avg:237.86ms |
| step:303/13210 train_time:72064ms step_avg:237.83ms |
| step:304/13210 train_time:72291ms step_avg:237.80ms |
| step:305/13210 train_time:72520ms step_avg:237.77ms |
| step:306/13210 train_time:72747ms step_avg:237.74ms |
| step:307/13210 train_time:72975ms step_avg:237.71ms |
| step:308/13210 train_time:73203ms step_avg:237.67ms |
| step:309/13210 train_time:73433ms step_avg:237.65ms |
| step:310/13210 train_time:73662ms step_avg:237.62ms |
| step:311/13210 train_time:73889ms step_avg:237.58ms |
| step:312/13210 train_time:74117ms step_avg:237.55ms |
| step:313/13210 train_time:74345ms step_avg:237.52ms |
| step:314/13210 train_time:74573ms step_avg:237.49ms |
| step:315/13210 train_time:74803ms step_avg:237.47ms |
| step:316/13210 train_time:75031ms step_avg:237.44ms |
| step:317/13210 train_time:75259ms step_avg:237.41ms |
| step:318/13210 train_time:75486ms step_avg:237.38ms |
| step:319/13210 train_time:75716ms step_avg:237.35ms |
| step:320/13210 train_time:75944ms step_avg:237.33ms |
| step:321/13210 train_time:76173ms step_avg:237.30ms |
| step:322/13210 train_time:76400ms step_avg:237.27ms |
| step:323/13210 train_time:76629ms step_avg:237.24ms |
| step:324/13210 train_time:76859ms step_avg:237.22ms |
| step:325/13210 train_time:77088ms step_avg:237.19ms |
| step:326/13210 train_time:77316ms step_avg:237.16ms |
| step:327/13210 train_time:77545ms step_avg:237.14ms |
| step:328/13210 train_time:77774ms step_avg:237.12ms |
| step:329/13210 train_time:78004ms step_avg:237.09ms |
| step:330/13210 train_time:78232ms step_avg:237.07ms |
| step:331/13210 train_time:78461ms step_avg:237.04ms |
| step:332/13210 train_time:78688ms step_avg:237.01ms |
| step:333/13210 train_time:78919ms step_avg:236.99ms |
| step:334/13210 train_time:79146ms step_avg:236.96ms |
| step:335/13210 train_time:79374ms step_avg:236.94ms |
| step:336/13210 train_time:79602ms step_avg:236.91ms |
| step:337/13210 train_time:79831ms step_avg:236.89ms |
| step:338/13210 train_time:80060ms step_avg:236.86ms |
| step:339/13210 train_time:80292ms step_avg:236.85ms |
| step:340/13210 train_time:80519ms step_avg:236.82ms |
| step:341/13210 train_time:80747ms step_avg:236.80ms |
| step:342/13210 train_time:80976ms step_avg:236.77ms |
| step:343/13210 train_time:81205ms step_avg:236.75ms |
| step:344/13210 train_time:81433ms step_avg:236.72ms |
| step:345/13210 train_time:81663ms step_avg:236.70ms |
| step:346/13210 train_time:81893ms step_avg:236.69ms |
| step:347/13210 train_time:82122ms step_avg:236.66ms |
| step:348/13210 train_time:82351ms step_avg:236.64ms |
| step:349/13210 train_time:82580ms step_avg:236.62ms |
| step:350/13210 train_time:82809ms step_avg:236.60ms |
| step:351/13210 train_time:83037ms step_avg:236.57ms |
| step:352/13210 train_time:83264ms step_avg:236.55ms |
| step:353/13210 train_time:83493ms step_avg:236.52ms |
| step:354/13210 train_time:83720ms step_avg:236.50ms |
| step:355/13210 train_time:83949ms step_avg:236.48ms |
| step:356/13210 train_time:84176ms step_avg:236.45ms |
| step:357/13210 train_time:84406ms step_avg:236.43ms |
| step:358/13210 train_time:84634ms step_avg:236.41ms |
| step:359/13210 train_time:84863ms step_avg:236.39ms |
| step:360/13210 train_time:85092ms step_avg:236.37ms |
| step:361/13210 train_time:85322ms step_avg:236.35ms |
| step:362/13210 train_time:85548ms step_avg:236.32ms |
| step:363/13210 train_time:85775ms step_avg:236.30ms |
| step:364/13210 train_time:86005ms step_avg:236.28ms |
| step:365/13210 train_time:86234ms step_avg:236.26ms |
| step:366/13210 train_time:86462ms step_avg:236.24ms |
| step:367/13210 train_time:86692ms step_avg:236.22ms |
| step:368/13210 train_time:86918ms step_avg:236.19ms |
| step:369/13210 train_time:87146ms step_avg:236.17ms |
| step:370/13210 train_time:87374ms step_avg:236.15ms |
| step:371/13210 train_time:87603ms step_avg:236.13ms |
| step:372/13210 train_time:87832ms step_avg:236.11ms |
| step:373/13210 train_time:88061ms step_avg:236.09ms |
| step:374/13210 train_time:88290ms step_avg:236.07ms |
| step:375/13210 train_time:88519ms step_avg:236.05ms |
| step:376/13210 train_time:88747ms step_avg:236.03ms |
| step:377/13210 train_time:88975ms step_avg:236.01ms |
| step:378/13210 train_time:89203ms step_avg:235.99ms |
| step:379/13210 train_time:89433ms step_avg:235.97ms |
| step:380/13210 train_time:89662ms step_avg:235.95ms |
| step:381/13210 train_time:89891ms step_avg:235.94ms |
| step:382/13210 train_time:90118ms step_avg:235.91ms |
| step:383/13210 train_time:90346ms step_avg:235.89ms |
| step:384/13210 train_time:90574ms step_avg:235.87ms |
| step:385/13210 train_time:90801ms step_avg:235.85ms |
| step:386/13210 train_time:91029ms step_avg:235.83ms |
| step:387/13210 train_time:91256ms step_avg:235.80ms |
| step:388/13210 train_time:91484ms step_avg:235.78ms |
| step:389/13210 train_time:91714ms step_avg:235.77ms |
| step:390/13210 train_time:91942ms step_avg:235.75ms |
| step:391/13210 train_time:92174ms step_avg:235.74ms |
| step:392/13210 train_time:92400ms step_avg:235.72ms |
| step:393/13210 train_time:92628ms step_avg:235.69ms |
| step:394/13210 train_time:92855ms step_avg:235.67ms |
| step:395/13210 train_time:93086ms step_avg:235.66ms |
| step:396/13210 train_time:93313ms step_avg:235.64ms |
| step:397/13210 train_time:93541ms step_avg:235.62ms |
| step:398/13210 train_time:93770ms step_avg:235.60ms |
| step:399/13210 train_time:93999ms step_avg:235.59ms |
| step:400/13210 train_time:94227ms step_avg:235.57ms |
| step:401/13210 train_time:94456ms step_avg:235.55ms |
| step:402/13210 train_time:94684ms step_avg:235.53ms |
| step:403/13210 train_time:94913ms step_avg:235.52ms |
| step:404/13210 train_time:95141ms step_avg:235.50ms |
| step:405/13210 train_time:95370ms step_avg:235.48ms |
| step:406/13210 train_time:95598ms step_avg:235.46ms |
| step:407/13210 train_time:95826ms step_avg:235.44ms |
| step:408/13210 train_time:96053ms step_avg:235.42ms |
| step:409/13210 train_time:96281ms step_avg:235.41ms |
| step:410/13210 train_time:96508ms step_avg:235.39ms |
| step:411/13210 train_time:96736ms step_avg:235.37ms |
| step:412/13210 train_time:96964ms step_avg:235.35ms |
| step:413/13210 train_time:97191ms step_avg:235.33ms |
| step:414/13210 train_time:97420ms step_avg:235.31ms |
| step:415/13210 train_time:97648ms step_avg:235.30ms |
| step:416/13210 train_time:97876ms step_avg:235.28ms |
| step:417/13210 train_time:98104ms step_avg:235.26ms |
| step:418/13210 train_time:98333ms step_avg:235.25ms |
| step:419/13210 train_time:98562ms step_avg:235.23ms |
| step:420/13210 train_time:98789ms step_avg:235.21ms |
| step:421/13210 train_time:99016ms step_avg:235.19ms |
| step:422/13210 train_time:99245ms step_avg:235.18ms |
| step:423/13210 train_time:99474ms step_avg:235.16ms |
| step:424/13210 train_time:99702ms step_avg:235.15ms |
| step:425/13210 train_time:99931ms step_avg:235.13ms |
| step:426/13210 train_time:100160ms step_avg:235.12ms |
| step:427/13210 train_time:100389ms step_avg:235.10ms |
| step:428/13210 train_time:100616ms step_avg:235.08ms |
| step:429/13210 train_time:100844ms step_avg:235.07ms |
| step:430/13210 train_time:101072ms step_avg:235.05ms |
| step:431/13210 train_time:101300ms step_avg:235.04ms |
| step:432/13210 train_time:101529ms step_avg:235.02ms |
| step:433/13210 train_time:101758ms step_avg:235.01ms |
| step:434/13210 train_time:101986ms step_avg:234.99ms |
| step:435/13210 train_time:102213ms step_avg:234.97ms |
| step:436/13210 train_time:102439ms step_avg:234.95ms |
| step:437/13210 train_time:102667ms step_avg:234.94ms |
| step:438/13210 train_time:102896ms step_avg:234.92ms |
| step:439/13210 train_time:103126ms step_avg:234.91ms |
| step:440/13210 train_time:103353ms step_avg:234.89ms |
| step:441/13210 train_time:103581ms step_avg:234.88ms |
| step:442/13210 train_time:103808ms step_avg:234.86ms |
| step:443/13210 train_time:104038ms step_avg:234.85ms |
| step:444/13210 train_time:104267ms step_avg:234.83ms |
| step:445/13210 train_time:104494ms step_avg:234.82ms |
| step:446/13210 train_time:104722ms step_avg:234.80ms |
| step:447/13210 train_time:104951ms step_avg:234.79ms |
| step:448/13210 train_time:105178ms step_avg:234.77ms |
| step:449/13210 train_time:105406ms step_avg:234.76ms |
| step:450/13210 train_time:105633ms step_avg:234.74ms |
| step:451/13210 train_time:105862ms step_avg:234.73ms |
| step:452/13210 train_time:106090ms step_avg:234.71ms |
| step:453/13210 train_time:106318ms step_avg:234.70ms |
| step:454/13210 train_time:106546ms step_avg:234.68ms |
| step:455/13210 train_time:106774ms step_avg:234.67ms |
| step:456/13210 train_time:107003ms step_avg:234.65ms |
| step:457/13210 train_time:107232ms step_avg:234.64ms |
| step:458/13210 train_time:107458ms step_avg:234.63ms |
| step:459/13210 train_time:107686ms step_avg:234.61ms |
| step:460/13210 train_time:107913ms step_avg:234.59ms |
| step:461/13210 train_time:108142ms step_avg:234.58ms |
| step:462/13210 train_time:108369ms step_avg:234.56ms |
| step:463/13210 train_time:108598ms step_avg:234.55ms |
| step:464/13210 train_time:108827ms step_avg:234.54ms |
| step:465/13210 train_time:109055ms step_avg:234.53ms |
| step:466/13210 train_time:109282ms step_avg:234.51ms |
| step:467/13210 train_time:109511ms step_avg:234.50ms |
| step:468/13210 train_time:109740ms step_avg:234.49ms |
| step:469/13210 train_time:109969ms step_avg:234.48ms |
| step:470/13210 train_time:110196ms step_avg:234.46ms |
| step:471/13210 train_time:110425ms step_avg:234.45ms |
| step:472/13210 train_time:110652ms step_avg:234.43ms |
| step:473/13210 train_time:110881ms step_avg:234.42ms |
| step:474/13210 train_time:111109ms step_avg:234.41ms |
| step:475/13210 train_time:111338ms step_avg:234.40ms |
| step:476/13210 train_time:111566ms step_avg:234.38ms |
| step:477/13210 train_time:111795ms step_avg:234.37ms |
| step:478/13210 train_time:112023ms step_avg:234.36ms |
| step:479/13210 train_time:112251ms step_avg:234.35ms |
| step:480/13210 train_time:112479ms step_avg:234.33ms |
| step:481/13210 train_time:112708ms step_avg:234.32ms |
| step:482/13210 train_time:112935ms step_avg:234.30ms |
| step:483/13210 train_time:113163ms step_avg:234.29ms |
| step:484/13210 train_time:113391ms step_avg:234.28ms |
| step:485/13210 train_time:113620ms step_avg:234.27ms |
| step:486/13210 train_time:113848ms step_avg:234.25ms |
| step:487/13210 train_time:114077ms step_avg:234.24ms |
| step:488/13210 train_time:114306ms step_avg:234.23ms |
| step:489/13210 train_time:114534ms step_avg:234.22ms |
| step:490/13210 train_time:114761ms step_avg:234.21ms |
| step:491/13210 train_time:114989ms step_avg:234.19ms |
| step:492/13210 train_time:115218ms step_avg:234.18ms |
| step:493/13210 train_time:115445ms step_avg:234.17ms |
| step:494/13210 train_time:115673ms step_avg:234.16ms |
| step:495/13210 train_time:115901ms step_avg:234.14ms |
| step:496/13210 train_time:116129ms step_avg:234.13ms |
| step:497/13210 train_time:116358ms step_avg:234.12ms |
| step:498/13210 train_time:116585ms step_avg:234.11ms |
| step:499/13210 train_time:116814ms step_avg:234.10ms |
| step:500/13210 train_time:117041ms step_avg:234.08ms |
| step:500/13210 val_loss:10.2386 train_time:117077ms step_avg:234.15ms |
| step:501/13210 train_time:121289ms step_avg:242.09ms |
| step:502/13210 train_time:121513ms step_avg:242.06ms |
| step:503/13210 train_time:121739ms step_avg:242.03ms |
| step:504/13210 train_time:121964ms step_avg:241.99ms |
| step:505/13210 train_time:122201ms step_avg:241.98ms |
| step:506/13210 train_time:122430ms step_avg:241.96ms |
| step:507/13210 train_time:122656ms step_avg:241.92ms |
| step:508/13210 train_time:122880ms step_avg:241.89ms |
| step:509/13210 train_time:123112ms step_avg:241.87ms |
| step:510/13210 train_time:123341ms step_avg:241.85ms |
| step:511/13210 train_time:123569ms step_avg:241.82ms |
| step:512/13210 train_time:123796ms step_avg:241.79ms |
| step:513/13210 train_time:124024ms step_avg:241.76ms |
| step:514/13210 train_time:124254ms step_avg:241.74ms |
| step:515/13210 train_time:124482ms step_avg:241.71ms |
| step:516/13210 train_time:124709ms step_avg:241.68ms |
| step:517/13210 train_time:124935ms step_avg:241.65ms |
| step:518/13210 train_time:125164ms step_avg:241.63ms |
| step:519/13210 train_time:125394ms step_avg:241.61ms |
| step:520/13210 train_time:125621ms step_avg:241.58ms |
| step:521/13210 train_time:125848ms step_avg:241.55ms |
| step:522/13210 train_time:126076ms step_avg:241.53ms |
| step:523/13210 train_time:126306ms step_avg:241.50ms |
| step:524/13210 train_time:126534ms step_avg:241.48ms |
| step:525/13210 train_time:126762ms step_avg:241.45ms |
| step:526/13210 train_time:126989ms step_avg:241.42ms |
| step:527/13210 train_time:127218ms step_avg:241.40ms |
| step:528/13210 train_time:127445ms step_avg:241.37ms |
| step:529/13210 train_time:127673ms step_avg:241.35ms |
| step:530/13210 train_time:127900ms step_avg:241.32ms |
| step:531/13210 train_time:128129ms step_avg:241.30ms |
| step:532/13210 train_time:128357ms step_avg:241.27ms |
| step:533/13210 train_time:128587ms step_avg:241.25ms |
| step:534/13210 train_time:128813ms step_avg:241.22ms |
| step:535/13210 train_time:129043ms step_avg:241.20ms |
| step:536/13210 train_time:129272ms step_avg:241.18ms |
| step:537/13210 train_time:129500ms step_avg:241.15ms |
| step:538/13210 train_time:129729ms step_avg:241.13ms |
| step:539/13210 train_time:129957ms step_avg:241.11ms |
| step:540/13210 train_time:130185ms step_avg:241.08ms |
| step:541/13210 train_time:130415ms step_avg:241.06ms |
| step:542/13210 train_time:130642ms step_avg:241.04ms |
| step:543/13210 train_time:130872ms step_avg:241.02ms |
| step:544/13210 train_time:131099ms step_avg:240.99ms |
| step:545/13210 train_time:131330ms step_avg:240.97ms |
| step:546/13210 train_time:131562ms step_avg:240.96ms |
| step:547/13210 train_time:131792ms step_avg:240.94ms |
| step:548/13210 train_time:132018ms step_avg:240.91ms |
| step:549/13210 train_time:132247ms step_avg:240.89ms |
| step:550/13210 train_time:132476ms step_avg:240.87ms |
| step:551/13210 train_time:132706ms step_avg:240.85ms |
| step:552/13210 train_time:132934ms step_avg:240.82ms |
| step:553/13210 train_time:133163ms step_avg:240.80ms |
| step:554/13210 train_time:133390ms step_avg:240.78ms |
| step:555/13210 train_time:133618ms step_avg:240.75ms |
| step:556/13210 train_time:133847ms step_avg:240.73ms |
| step:557/13210 train_time:134076ms step_avg:240.71ms |
| step:558/13210 train_time:134305ms step_avg:240.69ms |
| step:559/13210 train_time:134536ms step_avg:240.67ms |
| step:560/13210 train_time:134764ms step_avg:240.65ms |
| step:561/13210 train_time:134992ms step_avg:240.63ms |
| step:562/13210 train_time:135220ms step_avg:240.60ms |
| step:563/13210 train_time:135448ms step_avg:240.58ms |
| step:564/13210 train_time:135678ms step_avg:240.56ms |
| step:565/13210 train_time:135908ms step_avg:240.55ms |
| step:566/13210 train_time:136137ms step_avg:240.53ms |
| step:567/13210 train_time:136367ms step_avg:240.51ms |
| step:568/13210 train_time:136595ms step_avg:240.48ms |
| step:569/13210 train_time:136823ms step_avg:240.46ms |
| step:570/13210 train_time:137053ms step_avg:240.44ms |
| step:571/13210 train_time:137283ms step_avg:240.43ms |
| step:572/13210 train_time:137510ms step_avg:240.40ms |
| step:573/13210 train_time:137737ms step_avg:240.38ms |
| step:574/13210 train_time:137966ms step_avg:240.36ms |
| step:575/13210 train_time:138195ms step_avg:240.34ms |
| step:576/13210 train_time:138424ms step_avg:240.32ms |
| step:577/13210 train_time:138653ms step_avg:240.30ms |
| step:578/13210 train_time:138883ms step_avg:240.28ms |
| step:579/13210 train_time:139114ms step_avg:240.27ms |
| step:580/13210 train_time:139343ms step_avg:240.25ms |
| step:581/13210 train_time:139573ms step_avg:240.23ms |
| step:582/13210 train_time:139802ms step_avg:240.21ms |
| step:583/13210 train_time:140031ms step_avg:240.19ms |
| step:584/13210 train_time:140262ms step_avg:240.17ms |
| step:585/13210 train_time:140491ms step_avg:240.16ms |
| step:586/13210 train_time:140719ms step_avg:240.14ms |
| step:587/13210 train_time:140950ms step_avg:240.12ms |
| step:588/13210 train_time:141181ms step_avg:240.10ms |
| step:589/13210 train_time:141411ms step_avg:240.09ms |
| step:590/13210 train_time:141640ms step_avg:240.07ms |
| step:591/13210 train_time:141869ms step_avg:240.05ms |
| step:592/13210 train_time:142098ms step_avg:240.03ms |
| step:593/13210 train_time:142327ms step_avg:240.01ms |
| step:594/13210 train_time:142556ms step_avg:239.99ms |
| step:595/13210 train_time:142784ms step_avg:239.97ms |
| step:596/13210 train_time:143013ms step_avg:239.96ms |
| step:597/13210 train_time:143244ms step_avg:239.94ms |
| step:598/13210 train_time:143474ms step_avg:239.92ms |
| step:599/13210 train_time:143703ms step_avg:239.91ms |
| step:600/13210 train_time:143932ms step_avg:239.89ms |
| step:601/13210 train_time:144163ms step_avg:239.87ms |
| step:602/13210 train_time:144391ms step_avg:239.85ms |
| step:603/13210 train_time:144622ms step_avg:239.84ms |
| step:604/13210 train_time:144851ms step_avg:239.82ms |
| step:605/13210 train_time:145080ms step_avg:239.80ms |
| step:606/13210 train_time:145308ms step_avg:239.78ms |
| step:607/13210 train_time:145537ms step_avg:239.76ms |
| step:608/13210 train_time:145767ms step_avg:239.75ms |
| step:609/13210 train_time:145999ms step_avg:239.74ms |
| step:610/13210 train_time:146228ms step_avg:239.72ms |
| step:611/13210 train_time:146460ms step_avg:239.71ms |
| step:612/13210 train_time:146686ms step_avg:239.68ms |
| step:613/13210 train_time:146915ms step_avg:239.67ms |
| step:614/13210 train_time:147146ms step_avg:239.65ms |
| step:615/13210 train_time:147375ms step_avg:239.63ms |
| step:616/13210 train_time:147607ms step_avg:239.62ms |
| step:617/13210 train_time:147836ms step_avg:239.60ms |
| step:618/13210 train_time:148067ms step_avg:239.59ms |
| step:619/13210 train_time:148298ms step_avg:239.58ms |
| step:620/13210 train_time:148529ms step_avg:239.56ms |
| step:621/13210 train_time:148758ms step_avg:239.55ms |
| step:622/13210 train_time:148986ms step_avg:239.53ms |
| step:623/13210 train_time:149218ms step_avg:239.52ms |
| step:624/13210 train_time:149448ms step_avg:239.50ms |
| step:625/13210 train_time:149678ms step_avg:239.48ms |
| step:626/13210 train_time:149907ms step_avg:239.47ms |
| step:627/13210 train_time:150134ms step_avg:239.45ms |
| step:628/13210 train_time:150365ms step_avg:239.43ms |
| step:629/13210 train_time:150595ms step_avg:239.42ms |
| step:630/13210 train_time:150823ms step_avg:239.40ms |
| step:631/13210 train_time:151052ms step_avg:239.38ms |
| step:632/13210 train_time:151285ms step_avg:239.37ms |
| step:633/13210 train_time:151516ms step_avg:239.36ms |
| step:634/13210 train_time:151744ms step_avg:239.34ms |
| step:635/13210 train_time:151973ms step_avg:239.33ms |
| step:636/13210 train_time:152201ms step_avg:239.31ms |
| step:637/13210 train_time:152432ms step_avg:239.30ms |
| step:638/13210 train_time:152660ms step_avg:239.28ms |
| step:639/13210 train_time:152889ms step_avg:239.26ms |
| step:640/13210 train_time:153117ms step_avg:239.25ms |
| step:641/13210 train_time:153349ms step_avg:239.23ms |
| step:642/13210 train_time:153580ms step_avg:239.22ms |
| step:643/13210 train_time:153810ms step_avg:239.21ms |
| step:644/13210 train_time:154039ms step_avg:239.19ms |
| step:645/13210 train_time:154270ms step_avg:239.18ms |
| step:646/13210 train_time:154501ms step_avg:239.17ms |
| step:647/13210 train_time:154730ms step_avg:239.15ms |
| step:648/13210 train_time:154958ms step_avg:239.13ms |
| step:649/13210 train_time:155188ms step_avg:239.12ms |
| step:650/13210 train_time:155419ms step_avg:239.11ms |
| step:651/13210 train_time:155650ms step_avg:239.09ms |
| step:652/13210 train_time:155879ms step_avg:239.08ms |
| step:653/13210 train_time:156108ms step_avg:239.06ms |
| step:654/13210 train_time:156339ms step_avg:239.05ms |
| step:655/13210 train_time:156569ms step_avg:239.04ms |
| step:656/13210 train_time:156799ms step_avg:239.02ms |
| step:657/13210 train_time:157029ms step_avg:239.01ms |
| step:658/13210 train_time:157256ms step_avg:238.99ms |
| step:659/13210 train_time:157485ms step_avg:238.98ms |
| step:660/13210 train_time:157715ms step_avg:238.96ms |
| step:661/13210 train_time:157946ms step_avg:238.95ms |
| step:662/13210 train_time:158174ms step_avg:238.93ms |
| step:663/13210 train_time:158403ms step_avg:238.92ms |
| step:664/13210 train_time:158633ms step_avg:238.90ms |
| step:665/13210 train_time:158862ms step_avg:238.89ms |
| step:666/13210 train_time:159092ms step_avg:238.88ms |
| step:667/13210 train_time:159321ms step_avg:238.86ms |
| step:668/13210 train_time:159552ms step_avg:238.85ms |
| step:669/13210 train_time:159784ms step_avg:238.84ms |
| step:670/13210 train_time:160013ms step_avg:238.83ms |
| step:671/13210 train_time:160243ms step_avg:238.81ms |
| step:672/13210 train_time:160473ms step_avg:238.80ms |
| step:673/13210 train_time:160705ms step_avg:238.79ms |
| step:674/13210 train_time:160934ms step_avg:238.77ms |
| step:675/13210 train_time:161162ms step_avg:238.76ms |
| step:676/13210 train_time:161392ms step_avg:238.75ms |
| step:677/13210 train_time:161621ms step_avg:238.73ms |
| step:678/13210 train_time:161852ms step_avg:238.72ms |
| step:679/13210 train_time:162082ms step_avg:238.71ms |
| step:680/13210 train_time:162310ms step_avg:238.69ms |
| step:681/13210 train_time:162541ms step_avg:238.68ms |
| step:682/13210 train_time:162770ms step_avg:238.67ms |
| step:683/13210 train_time:163000ms step_avg:238.65ms |
| step:684/13210 train_time:163231ms step_avg:238.64ms |
| step:685/13210 train_time:163461ms step_avg:238.63ms |
| step:686/13210 train_time:163690ms step_avg:238.62ms |
| step:687/13210 train_time:163921ms step_avg:238.60ms |
| step:688/13210 train_time:164149ms step_avg:238.59ms |
| step:689/13210 train_time:164379ms step_avg:238.58ms |
| step:690/13210 train_time:164607ms step_avg:238.56ms |
| step:691/13210 train_time:164838ms step_avg:238.55ms |
| step:692/13210 train_time:165065ms step_avg:238.53ms |
| step:693/13210 train_time:165295ms step_avg:238.52ms |
| step:694/13210 train_time:165523ms step_avg:238.51ms |
| step:695/13210 train_time:165751ms step_avg:238.49ms |
| step:696/13210 train_time:165980ms step_avg:238.48ms |
| step:697/13210 train_time:166212ms step_avg:238.47ms |
| step:698/13210 train_time:166441ms step_avg:238.45ms |
| step:699/13210 train_time:166673ms step_avg:238.45ms |
| step:700/13210 train_time:166903ms step_avg:238.43ms |
| step:701/13210 train_time:167132ms step_avg:238.42ms |
| step:702/13210 train_time:167362ms step_avg:238.41ms |
| step:703/13210 train_time:167592ms step_avg:238.40ms |
| step:704/13210 train_time:167821ms step_avg:238.38ms |
| step:705/13210 train_time:168051ms step_avg:238.37ms |
| step:706/13210 train_time:168279ms step_avg:238.35ms |
| step:707/13210 train_time:168509ms step_avg:238.34ms |
| step:708/13210 train_time:168738ms step_avg:238.33ms |
| step:709/13210 train_time:168968ms step_avg:238.32ms |
| step:710/13210 train_time:169196ms step_avg:238.30ms |
| step:711/13210 train_time:169425ms step_avg:238.29ms |
| step:712/13210 train_time:169654ms step_avg:238.28ms |
| step:713/13210 train_time:169883ms step_avg:238.26ms |
| step:714/13210 train_time:170113ms step_avg:238.25ms |
| step:715/13210 train_time:170343ms step_avg:238.24ms |
| step:716/13210 train_time:170571ms step_avg:238.23ms |
| step:717/13210 train_time:170801ms step_avg:238.22ms |
| step:718/13210 train_time:171030ms step_avg:238.20ms |
| step:719/13210 train_time:171261ms step_avg:238.19ms |
| step:720/13210 train_time:171490ms step_avg:238.18ms |
| step:721/13210 train_time:171721ms step_avg:238.17ms |
| step:722/13210 train_time:171949ms step_avg:238.16ms |
| step:723/13210 train_time:172178ms step_avg:238.14ms |
| step:724/13210 train_time:172406ms step_avg:238.13ms |
| step:725/13210 train_time:172636ms step_avg:238.12ms |
| step:726/13210 train_time:172864ms step_avg:238.11ms |
| step:727/13210 train_time:173094ms step_avg:238.09ms |
| step:728/13210 train_time:173322ms step_avg:238.08ms |
| step:729/13210 train_time:173551ms step_avg:238.07ms |
| step:730/13210 train_time:173778ms step_avg:238.05ms |
| step:731/13210 train_time:174006ms step_avg:238.04ms |
| step:732/13210 train_time:174234ms step_avg:238.03ms |
| step:733/13210 train_time:174463ms step_avg:238.01ms |
| step:734/13210 train_time:174691ms step_avg:238.00ms |
| step:735/13210 train_time:174920ms step_avg:237.99ms |
| step:736/13210 train_time:175149ms step_avg:237.97ms |
| step:737/13210 train_time:175379ms step_avg:237.96ms |
| step:738/13210 train_time:175608ms step_avg:237.95ms |
| step:739/13210 train_time:175836ms step_avg:237.94ms |
| step:740/13210 train_time:176065ms step_avg:237.93ms |
| step:741/13210 train_time:176295ms step_avg:237.91ms |
| step:742/13210 train_time:176524ms step_avg:237.90ms |
| step:743/13210 train_time:176754ms step_avg:237.89ms |
| step:744/13210 train_time:176981ms step_avg:237.88ms |
| step:745/13210 train_time:177211ms step_avg:237.87ms |
| step:746/13210 train_time:177441ms step_avg:237.86ms |
| step:747/13210 train_time:177669ms step_avg:237.84ms |
| step:748/13210 train_time:177897ms step_avg:237.83ms |
| step:749/13210 train_time:178127ms step_avg:237.82ms |
| step:750/13210 train_time:178356ms step_avg:237.81ms |
| step:750/13210 val_loss:10.2148 train_time:178392ms step_avg:237.86ms |
| step:751/13210 train_time:178589ms step_avg:237.80ms |
| step:752/13210 train_time:178816ms step_avg:237.79ms |
| step:753/13210 train_time:179047ms step_avg:237.78ms |
| step:754/13210 train_time:179273ms step_avg:237.76ms |
| step:755/13210 train_time:179501ms step_avg:237.75ms |
| step:756/13210 train_time:179730ms step_avg:237.74ms |
| step:757/13210 train_time:179959ms step_avg:237.73ms |
| step:758/13210 train_time:180186ms step_avg:237.71ms |
| step:759/13210 train_time:180415ms step_avg:237.70ms |
| step:760/13210 train_time:180646ms step_avg:237.69ms |
| step:761/13210 train_time:180875ms step_avg:237.68ms |
| step:762/13210 train_time:181104ms step_avg:237.67ms |
| step:763/13210 train_time:181334ms step_avg:237.66ms |
| step:764/13210 train_time:181562ms step_avg:237.65ms |
| step:765/13210 train_time:181793ms step_avg:237.64ms |
| step:766/13210 train_time:182022ms step_avg:237.63ms |
| step:767/13210 train_time:182250ms step_avg:237.61ms |
| step:768/13210 train_time:182478ms step_avg:237.60ms |
| step:769/13210 train_time:182707ms step_avg:237.59ms |
| step:770/13210 train_time:182938ms step_avg:237.58ms |
| step:771/13210 train_time:183168ms step_avg:237.57ms |
| step:772/13210 train_time:183396ms step_avg:237.56ms |
| step:773/13210 train_time:183624ms step_avg:237.55ms |
| step:774/13210 train_time:183852ms step_avg:237.54ms |
| step:775/13210 train_time:184080ms step_avg:237.52ms |
| step:776/13210 train_time:184309ms step_avg:237.51ms |
| step:777/13210 train_time:184536ms step_avg:237.50ms |
| step:778/13210 train_time:184764ms step_avg:237.49ms |
| step:779/13210 train_time:184994ms step_avg:237.48ms |
| step:780/13210 train_time:185222ms step_avg:237.46ms |
| step:781/13210 train_time:185450ms step_avg:237.45ms |
| step:782/13210 train_time:185680ms step_avg:237.44ms |
| step:783/13210 train_time:185909ms step_avg:237.43ms |
| step:784/13210 train_time:186137ms step_avg:237.42ms |
| step:785/13210 train_time:186366ms step_avg:237.41ms |
| step:786/13210 train_time:186594ms step_avg:237.40ms |
| step:787/13210 train_time:186822ms step_avg:237.39ms |
| step:788/13210 train_time:187051ms step_avg:237.37ms |
| step:789/13210 train_time:187279ms step_avg:237.36ms |
| step:790/13210 train_time:187508ms step_avg:237.35ms |
| step:791/13210 train_time:187738ms step_avg:237.34ms |
| step:792/13210 train_time:187965ms step_avg:237.33ms |
| step:793/13210 train_time:188193ms step_avg:237.32ms |
| step:794/13210 train_time:188422ms step_avg:237.31ms |
| step:795/13210 train_time:188650ms step_avg:237.30ms |
| step:796/13210 train_time:188879ms step_avg:237.29ms |
| step:797/13210 train_time:189110ms step_avg:237.28ms |
| step:798/13210 train_time:189337ms step_avg:237.26ms |
| step:799/13210 train_time:189566ms step_avg:237.25ms |
| step:800/13210 train_time:189795ms step_avg:237.24ms |
| step:801/13210 train_time:190025ms step_avg:237.23ms |
| step:802/13210 train_time:190254ms step_avg:237.22ms |
| step:803/13210 train_time:190486ms step_avg:237.22ms |
| step:804/13210 train_time:190713ms step_avg:237.20ms |
| step:805/13210 train_time:190941ms step_avg:237.19ms |
| step:806/13210 train_time:191168ms step_avg:237.18ms |
| step:807/13210 train_time:191398ms step_avg:237.17ms |
| step:808/13210 train_time:191628ms step_avg:237.16ms |
| step:809/13210 train_time:191857ms step_avg:237.15ms |
| step:810/13210 train_time:192087ms step_avg:237.14ms |
| step:811/13210 train_time:192316ms step_avg:237.13ms |
| step:812/13210 train_time:192545ms step_avg:237.12ms |
| step:813/13210 train_time:192775ms step_avg:237.12ms |
| step:814/13210 train_time:193004ms step_avg:237.11ms |
| step:815/13210 train_time:193230ms step_avg:237.09ms |
| step:816/13210 train_time:193459ms step_avg:237.08ms |
| step:817/13210 train_time:193690ms step_avg:237.07ms |
| step:818/13210 train_time:193919ms step_avg:237.06ms |
| step:819/13210 train_time:194146ms step_avg:237.05ms |
| step:820/13210 train_time:194375ms step_avg:237.04ms |
| step:821/13210 train_time:194606ms step_avg:237.03ms |
| step:822/13210 train_time:194832ms step_avg:237.02ms |
| step:823/13210 train_time:195062ms step_avg:237.01ms |
| step:824/13210 train_time:195289ms step_avg:237.00ms |
| step:825/13210 train_time:195519ms step_avg:236.99ms |
| step:826/13210 train_time:195747ms step_avg:236.98ms |
| step:827/13210 train_time:195976ms step_avg:236.97ms |
| step:828/13210 train_time:196203ms step_avg:236.96ms |
| step:829/13210 train_time:196433ms step_avg:236.95ms |
| step:830/13210 train_time:196661ms step_avg:236.94ms |
| step:831/13210 train_time:196889ms step_avg:236.93ms |
| step:832/13210 train_time:197118ms step_avg:236.92ms |
| step:833/13210 train_time:197346ms step_avg:236.91ms |
| step:834/13210 train_time:197577ms step_avg:236.90ms |
| step:835/13210 train_time:197807ms step_avg:236.89ms |
| step:836/13210 train_time:198035ms step_avg:236.88ms |
| step:837/13210 train_time:198263ms step_avg:236.87ms |
| step:838/13210 train_time:198491ms step_avg:236.86ms |
| step:839/13210 train_time:198720ms step_avg:236.85ms |
| step:840/13210 train_time:198949ms step_avg:236.84ms |
| step:841/13210 train_time:199178ms step_avg:236.83ms |
| step:842/13210 train_time:199405ms step_avg:236.82ms |
| step:843/13210 train_time:199632ms step_avg:236.81ms |
| step:844/13210 train_time:199861ms step_avg:236.80ms |
| step:845/13210 train_time:200090ms step_avg:236.79ms |
| step:846/13210 train_time:200317ms step_avg:236.78ms |
| step:847/13210 train_time:200546ms step_avg:236.77ms |
| step:848/13210 train_time:200776ms step_avg:236.76ms |
| step:849/13210 train_time:201005ms step_avg:236.76ms |
| step:850/13210 train_time:201234ms step_avg:236.75ms |
| step:851/13210 train_time:201464ms step_avg:236.74ms |
| step:852/13210 train_time:201695ms step_avg:236.73ms |
| step:853/13210 train_time:201923ms step_avg:236.72ms |
| step:854/13210 train_time:202151ms step_avg:236.71ms |
| step:855/13210 train_time:202381ms step_avg:236.70ms |
| step:856/13210 train_time:202610ms step_avg:236.69ms |
| step:857/13210 train_time:202838ms step_avg:236.68ms |
| step:858/13210 train_time:203067ms step_avg:236.67ms |
| step:859/13210 train_time:203295ms step_avg:236.66ms |
| step:860/13210 train_time:203524ms step_avg:236.66ms |
| step:861/13210 train_time:203755ms step_avg:236.65ms |
| step:862/13210 train_time:203982ms step_avg:236.64ms |
| step:863/13210 train_time:204211ms step_avg:236.63ms |
| step:864/13210 train_time:204440ms step_avg:236.62ms |
| step:865/13210 train_time:204667ms step_avg:236.61ms |
| step:866/13210 train_time:204895ms step_avg:236.60ms |
| step:867/13210 train_time:205123ms step_avg:236.59ms |
| step:868/13210 train_time:205350ms step_avg:236.58ms |
| step:869/13210 train_time:205580ms step_avg:236.57ms |
| step:870/13210 train_time:205809ms step_avg:236.56ms |
| step:871/13210 train_time:206038ms step_avg:236.55ms |
| step:872/13210 train_time:206266ms step_avg:236.54ms |
| step:873/13210 train_time:206496ms step_avg:236.54ms |
| step:874/13210 train_time:206723ms step_avg:236.52ms |
| step:875/13210 train_time:206952ms step_avg:236.52ms |
| step:876/13210 train_time:207181ms step_avg:236.51ms |
| step:877/13210 train_time:207410ms step_avg:236.50ms |
| step:878/13210 train_time:207637ms step_avg:236.49ms |
| step:879/13210 train_time:207866ms step_avg:236.48ms |
| step:880/13210 train_time:208094ms step_avg:236.47ms |
| step:881/13210 train_time:208324ms step_avg:236.46ms |
| step:882/13210 train_time:208551ms step_avg:236.45ms |
| step:883/13210 train_time:208780ms step_avg:236.44ms |
| step:884/13210 train_time:209008ms step_avg:236.43ms |
| step:885/13210 train_time:209238ms step_avg:236.43ms |
| step:886/13210 train_time:209465ms step_avg:236.42ms |
| step:887/13210 train_time:209696ms step_avg:236.41ms |
| step:888/13210 train_time:209924ms step_avg:236.40ms |
| step:889/13210 train_time:210152ms step_avg:236.39ms |
| step:890/13210 train_time:210382ms step_avg:236.38ms |
| step:891/13210 train_time:210611ms step_avg:236.38ms |
| step:892/13210 train_time:210838ms step_avg:236.37ms |
| step:893/13210 train_time:211067ms step_avg:236.36ms |
| step:894/13210 train_time:211296ms step_avg:236.35ms |
| step:895/13210 train_time:211525ms step_avg:236.34ms |
| step:896/13210 train_time:211754ms step_avg:236.33ms |
| step:897/13210 train_time:211982ms step_avg:236.32ms |
| step:898/13210 train_time:212210ms step_avg:236.31ms |
| step:899/13210 train_time:212438ms step_avg:236.31ms |
| step:900/13210 train_time:212668ms step_avg:236.30ms |
| step:901/13210 train_time:212898ms step_avg:236.29ms |
| step:902/13210 train_time:213125ms step_avg:236.28ms |
| step:903/13210 train_time:213354ms step_avg:236.27ms |
| step:904/13210 train_time:213585ms step_avg:236.27ms |
| step:905/13210 train_time:213815ms step_avg:236.26ms |
| step:906/13210 train_time:214043ms step_avg:236.25ms |
| step:907/13210 train_time:214273ms step_avg:236.24ms |
| step:908/13210 train_time:214500ms step_avg:236.23ms |
| step:909/13210 train_time:214729ms step_avg:236.23ms |
| step:910/13210 train_time:214958ms step_avg:236.22ms |
| step:911/13210 train_time:215187ms step_avg:236.21ms |
| step:912/13210 train_time:215416ms step_avg:236.20ms |
| step:913/13210 train_time:215646ms step_avg:236.19ms |
| step:914/13210 train_time:215874ms step_avg:236.19ms |
| step:915/13210 train_time:216102ms step_avg:236.18ms |
| step:916/13210 train_time:216330ms step_avg:236.17ms |
| step:917/13210 train_time:216560ms step_avg:236.16ms |
| step:918/13210 train_time:216789ms step_avg:236.15ms |
| step:919/13210 train_time:217016ms step_avg:236.14ms |
| step:920/13210 train_time:217245ms step_avg:236.14ms |
| step:921/13210 train_time:217475ms step_avg:236.13ms |
| step:922/13210 train_time:217702ms step_avg:236.12ms |
| step:923/13210 train_time:217931ms step_avg:236.11ms |
| step:924/13210 train_time:218161ms step_avg:236.10ms |
| step:925/13210 train_time:218389ms step_avg:236.10ms |
| step:926/13210 train_time:218617ms step_avg:236.09ms |
| step:927/13210 train_time:218845ms step_avg:236.08ms |
| step:928/13210 train_time:219074ms step_avg:236.07ms |
| step:929/13210 train_time:219304ms step_avg:236.06ms |
| step:930/13210 train_time:219532ms step_avg:236.06ms |
| step:931/13210 train_time:219761ms step_avg:236.05ms |
| step:932/13210 train_time:219987ms step_avg:236.04ms |
| step:933/13210 train_time:220218ms step_avg:236.03ms |
| step:934/13210 train_time:220447ms step_avg:236.02ms |
| step:935/13210 train_time:220675ms step_avg:236.02ms |
| step:936/13210 train_time:220903ms step_avg:236.01ms |
| step:937/13210 train_time:221132ms step_avg:236.00ms |
| step:938/13210 train_time:221362ms step_avg:235.99ms |
| step:939/13210 train_time:221594ms step_avg:235.99ms |
| step:940/13210 train_time:221823ms step_avg:235.98ms |
| step:941/13210 train_time:222050ms step_avg:235.97ms |
| step:942/13210 train_time:222282ms step_avg:235.97ms |
| step:943/13210 train_time:222511ms step_avg:235.96ms |
| step:944/13210 train_time:222738ms step_avg:235.95ms |
| step:945/13210 train_time:222967ms step_avg:235.94ms |
| step:946/13210 train_time:223195ms step_avg:235.94ms |
| step:947/13210 train_time:223426ms step_avg:235.93ms |
| step:948/13210 train_time:223654ms step_avg:235.92ms |
| step:949/13210 train_time:223883ms step_avg:235.91ms |
| step:950/13210 train_time:224111ms step_avg:235.91ms |
| step:951/13210 train_time:224341ms step_avg:235.90ms |
| step:952/13210 train_time:224570ms step_avg:235.89ms |
| step:953/13210 train_time:224799ms step_avg:235.89ms |
| step:954/13210 train_time:225027ms step_avg:235.88ms |
| step:955/13210 train_time:225256ms step_avg:235.87ms |
| step:956/13210 train_time:225486ms step_avg:235.86ms |
| step:957/13210 train_time:225716ms step_avg:235.86ms |
| step:958/13210 train_time:225945ms step_avg:235.85ms |
| step:959/13210 train_time:226174ms step_avg:235.84ms |
| step:960/13210 train_time:226403ms step_avg:235.84ms |
| step:961/13210 train_time:226632ms step_avg:235.83ms |
| step:962/13210 train_time:226860ms step_avg:235.82ms |
| step:963/13210 train_time:227088ms step_avg:235.81ms |
| step:964/13210 train_time:227317ms step_avg:235.81ms |
| step:965/13210 train_time:227546ms step_avg:235.80ms |
| step:966/13210 train_time:227775ms step_avg:235.79ms |
| step:967/13210 train_time:228003ms step_avg:235.78ms |
| step:968/13210 train_time:228230ms step_avg:235.78ms |
| step:969/13210 train_time:228458ms step_avg:235.77ms |
| step:970/13210 train_time:228685ms step_avg:235.76ms |
| step:971/13210 train_time:228916ms step_avg:235.75ms |
| step:972/13210 train_time:229145ms step_avg:235.75ms |
| step:973/13210 train_time:229375ms step_avg:235.74ms |
| step:974/13210 train_time:229602ms step_avg:235.73ms |
| step:975/13210 train_time:229831ms step_avg:235.72ms |
| step:976/13210 train_time:230058ms step_avg:235.72ms |
| step:977/13210 train_time:230289ms step_avg:235.71ms |
| step:978/13210 train_time:230518ms step_avg:235.70ms |
| step:979/13210 train_time:230747ms step_avg:235.70ms |
| step:980/13210 train_time:230974ms step_avg:235.69ms |
| step:981/13210 train_time:231201ms step_avg:235.68ms |
| step:982/13210 train_time:231429ms step_avg:235.67ms |
| step:983/13210 train_time:231657ms step_avg:235.66ms |
| step:984/13210 train_time:231887ms step_avg:235.66ms |
| step:985/13210 train_time:232115ms step_avg:235.65ms |
| step:986/13210 train_time:232345ms step_avg:235.64ms |
| step:987/13210 train_time:232574ms step_avg:235.64ms |
| step:988/13210 train_time:232802ms step_avg:235.63ms |
| step:989/13210 train_time:233030ms step_avg:235.62ms |
| step:990/13210 train_time:233260ms step_avg:235.62ms |
| step:991/13210 train_time:233489ms step_avg:235.61ms |
| step:992/13210 train_time:233717ms step_avg:235.60ms |
| step:993/13210 train_time:233945ms step_avg:235.59ms |
| step:994/13210 train_time:234173ms step_avg:235.59ms |
| step:995/13210 train_time:234402ms step_avg:235.58ms |
| step:996/13210 train_time:234630ms step_avg:235.57ms |
| step:997/13210 train_time:234859ms step_avg:235.57ms |
| step:998/13210 train_time:235087ms step_avg:235.56ms |
| step:999/13210 train_time:235317ms step_avg:235.55ms |
| step:1000/13210 train_time:235549ms step_avg:235.55ms |
| step:1000/13210 val_loss:10.2397 train_time:235585ms step_avg:235.59ms |
| step:1001/13210 train_time:239784ms step_avg:239.54ms |
| step:1002/13210 train_time:240008ms step_avg:239.53ms |
| step:1003/13210 train_time:240234ms step_avg:239.52ms |
| step:1004/13210 train_time:240460ms step_avg:239.50ms |
| step:1005/13210 train_time:240696ms step_avg:239.50ms |
| step:1006/13210 train_time:240923ms step_avg:239.49ms |
| step:1007/13210 train_time:241150ms step_avg:239.47ms |
| step:1008/13210 train_time:241374ms step_avg:239.46ms |
| step:1009/13210 train_time:241606ms step_avg:239.45ms |
| step:1010/13210 train_time:241834ms step_avg:239.44ms |
| step:1011/13210 train_time:242062ms step_avg:239.43ms |
| step:1012/13210 train_time:242288ms step_avg:239.41ms |
| step:1013/13210 train_time:242516ms step_avg:239.40ms |
| step:1014/13210 train_time:242747ms step_avg:239.40ms |
| step:1015/13210 train_time:242976ms step_avg:239.38ms |
| step:1016/13210 train_time:243201ms step_avg:239.37ms |
| step:1017/13210 train_time:243429ms step_avg:239.36ms |
| step:1018/13210 train_time:243657ms step_avg:239.35ms |
| step:1019/13210 train_time:243887ms step_avg:239.34ms |
| step:1020/13210 train_time:244114ms step_avg:239.33ms |
| step:1021/13210 train_time:244341ms step_avg:239.32ms |
| step:1022/13210 train_time:244570ms step_avg:239.31ms |
| step:1023/13210 train_time:244798ms step_avg:239.29ms |
| step:1024/13210 train_time:245026ms step_avg:239.28ms |
| step:1025/13210 train_time:245254ms step_avg:239.27ms |
| step:1026/13210 train_time:245480ms step_avg:239.26ms |
| step:1027/13210 train_time:245708ms step_avg:239.25ms |
| step:1028/13210 train_time:245936ms step_avg:239.24ms |
| step:1029/13210 train_time:246163ms step_avg:239.23ms |
| step:1030/13210 train_time:246391ms step_avg:239.21ms |
| step:1031/13210 train_time:246617ms step_avg:239.20ms |
| step:1032/13210 train_time:246847ms step_avg:239.19ms |
| step:1033/13210 train_time:247076ms step_avg:239.18ms |
| step:1034/13210 train_time:247305ms step_avg:239.17ms |
| step:1035/13210 train_time:247533ms step_avg:239.16ms |
| step:1036/13210 train_time:247762ms step_avg:239.15ms |
| step:1037/13210 train_time:247993ms step_avg:239.14ms |
| step:1038/13210 train_time:248220ms step_avg:239.13ms |
| step:1039/13210 train_time:248449ms step_avg:239.12ms |
| step:1040/13210 train_time:248677ms step_avg:239.11ms |
| step:1041/13210 train_time:248907ms step_avg:239.10ms |
| step:1042/13210 train_time:249137ms step_avg:239.10ms |
| step:1043/13210 train_time:249365ms step_avg:239.08ms |
| step:1044/13210 train_time:249592ms step_avg:239.07ms |
| step:1045/13210 train_time:249821ms step_avg:239.06ms |
| step:1046/13210 train_time:250049ms step_avg:239.05ms |
| step:1047/13210 train_time:250277ms step_avg:239.04ms |
| step:1048/13210 train_time:250504ms step_avg:239.03ms |
| step:1049/13210 train_time:250734ms step_avg:239.02ms |
| step:1050/13210 train_time:250964ms step_avg:239.01ms |
| step:1051/13210 train_time:251193ms step_avg:239.00ms |
| step:1052/13210 train_time:251420ms step_avg:238.99ms |
| step:1053/13210 train_time:251651ms step_avg:238.99ms |
| step:1054/13210 train_time:251879ms step_avg:238.97ms |
| step:1055/13210 train_time:252108ms step_avg:238.96ms |
| step:1056/13210 train_time:252337ms step_avg:238.96ms |
| step:1057/13210 train_time:252566ms step_avg:238.95ms |
| step:1058/13210 train_time:252796ms step_avg:238.94ms |
| step:1059/13210 train_time:253026ms step_avg:238.93ms |
| step:1060/13210 train_time:253256ms step_avg:238.92ms |
| step:1061/13210 train_time:253485ms step_avg:238.91ms |
| step:1062/13210 train_time:253714ms step_avg:238.90ms |
| step:1063/13210 train_time:253944ms step_avg:238.89ms |
| step:1064/13210 train_time:254171ms step_avg:238.88ms |
| step:1065/13210 train_time:254402ms step_avg:238.88ms |
| step:1066/13210 train_time:254629ms step_avg:238.86ms |
| step:1067/13210 train_time:254858ms step_avg:238.85ms |
| step:1068/13210 train_time:255084ms step_avg:238.84ms |
| step:1069/13210 train_time:255313ms step_avg:238.83ms |
| step:1070/13210 train_time:255541ms step_avg:238.82ms |
| step:1071/13210 train_time:255771ms step_avg:238.82ms |
| step:1072/13210 train_time:255999ms step_avg:238.80ms |
| step:1073/13210 train_time:256230ms step_avg:238.80ms |
| step:1074/13210 train_time:256458ms step_avg:238.79ms |
| step:1075/13210 train_time:256687ms step_avg:238.78ms |
| step:1076/13210 train_time:256915ms step_avg:238.77ms |
| step:1077/13210 train_time:257144ms step_avg:238.76ms |
| step:1078/13210 train_time:257372ms step_avg:238.75ms |
| step:1079/13210 train_time:257603ms step_avg:238.74ms |
| step:1080/13210 train_time:257831ms step_avg:238.73ms |
| step:1081/13210 train_time:258060ms step_avg:238.72ms |
| step:1082/13210 train_time:258289ms step_avg:238.71ms |
| step:1083/13210 train_time:258521ms step_avg:238.71ms |
| step:1084/13210 train_time:258748ms step_avg:238.70ms |
| step:1085/13210 train_time:258976ms step_avg:238.69ms |
| step:1086/13210 train_time:259204ms step_avg:238.68ms |
| step:1087/13210 train_time:259435ms step_avg:238.67ms |
| step:1088/13210 train_time:259667ms step_avg:238.66ms |
| step:1089/13210 train_time:259896ms step_avg:238.66ms |
| step:1090/13210 train_time:260125ms step_avg:238.65ms |
| step:1091/13210 train_time:260352ms step_avg:238.64ms |
| step:1092/13210 train_time:260580ms step_avg:238.63ms |
| step:1093/13210 train_time:260809ms step_avg:238.62ms |
| step:1094/13210 train_time:261038ms step_avg:238.61ms |
| step:1095/13210 train_time:261269ms step_avg:238.60ms |
| step:1096/13210 train_time:261496ms step_avg:238.59ms |
| step:1097/13210 train_time:261723ms step_avg:238.58ms |
| step:1098/13210 train_time:261950ms step_avg:238.57ms |
| step:1099/13210 train_time:262180ms step_avg:238.56ms |
| step:1100/13210 train_time:262408ms step_avg:238.55ms |
| step:1101/13210 train_time:262638ms step_avg:238.55ms |
| step:1102/13210 train_time:262867ms step_avg:238.54ms |
| step:1103/13210 train_time:263096ms step_avg:238.53ms |
| step:1104/13210 train_time:263328ms step_avg:238.52ms |
| step:1105/13210 train_time:263556ms step_avg:238.51ms |
| step:1106/13210 train_time:263785ms step_avg:238.50ms |
| step:1107/13210 train_time:264015ms step_avg:238.50ms |
| step:1108/13210 train_time:264243ms step_avg:238.49ms |
| step:1109/13210 train_time:264475ms step_avg:238.48ms |
| step:1110/13210 train_time:264704ms step_avg:238.47ms |
| step:1111/13210 train_time:264934ms step_avg:238.46ms |
| step:1112/13210 train_time:265161ms step_avg:238.45ms |
| step:1113/13210 train_time:265394ms step_avg:238.45ms |
| step:1114/13210 train_time:265623ms step_avg:238.44ms |
| step:1115/13210 train_time:265852ms step_avg:238.43ms |
| step:1116/13210 train_time:266081ms step_avg:238.42ms |
| step:1117/13210 train_time:266310ms step_avg:238.42ms |
| step:1118/13210 train_time:266538ms step_avg:238.41ms |
| step:1119/13210 train_time:266768ms step_avg:238.40ms |
| step:1120/13210 train_time:266995ms step_avg:238.39ms |
| step:1121/13210 train_time:267224ms step_avg:238.38ms |
| step:1122/13210 train_time:267454ms step_avg:238.37ms |
| step:1123/13210 train_time:267684ms step_avg:238.36ms |
| step:1124/13210 train_time:267912ms step_avg:238.36ms |
| step:1125/13210 train_time:268140ms step_avg:238.35ms |
| step:1126/13210 train_time:268370ms step_avg:238.34ms |
| step:1127/13210 train_time:268600ms step_avg:238.33ms |
| step:1128/13210 train_time:268828ms step_avg:238.32ms |
| step:1129/13210 train_time:269058ms step_avg:238.32ms |
| step:1130/13210 train_time:269286ms step_avg:238.31ms |
| step:1131/13210 train_time:269515ms step_avg:238.30ms |
| step:1132/13210 train_time:269744ms step_avg:238.29ms |
| step:1133/13210 train_time:269972ms step_avg:238.28ms |
| step:1134/13210 train_time:270199ms step_avg:238.27ms |
| step:1135/13210 train_time:270428ms step_avg:238.26ms |
| step:1136/13210 train_time:270656ms step_avg:238.25ms |
| step:1137/13210 train_time:270885ms step_avg:238.25ms |
| step:1138/13210 train_time:271111ms step_avg:238.24ms |
| step:1139/13210 train_time:271339ms step_avg:238.23ms |
| step:1140/13210 train_time:271568ms step_avg:238.22ms |
| step:1141/13210 train_time:271798ms step_avg:238.21ms |
| step:1142/13210 train_time:272027ms step_avg:238.20ms |
| step:1143/13210 train_time:272255ms step_avg:238.19ms |
| step:1144/13210 train_time:272483ms step_avg:238.18ms |
| step:1145/13210 train_time:272710ms step_avg:238.17ms |
| step:1146/13210 train_time:272940ms step_avg:238.17ms |
| step:1147/13210 train_time:273169ms step_avg:238.16ms |
| step:1148/13210 train_time:273396ms step_avg:238.15ms |
| step:1149/13210 train_time:273627ms step_avg:238.14ms |
| step:1150/13210 train_time:273854ms step_avg:238.13ms |
| step:1151/13210 train_time:274085ms step_avg:238.13ms |
| step:1152/13210 train_time:274312ms step_avg:238.12ms |
| step:1153/13210 train_time:274544ms step_avg:238.11ms |
| step:1154/13210 train_time:274771ms step_avg:238.10ms |
| step:1155/13210 train_time:275001ms step_avg:238.10ms |
| step:1156/13210 train_time:275229ms step_avg:238.09ms |
| step:1157/13210 train_time:275459ms step_avg:238.08ms |
| step:1158/13210 train_time:275688ms step_avg:238.07ms |
| step:1159/13210 train_time:275919ms step_avg:238.07ms |
| step:1160/13210 train_time:276146ms step_avg:238.06ms |
| step:1161/13210 train_time:276375ms step_avg:238.05ms |
| step:1162/13210 train_time:276603ms step_avg:238.04ms |
| step:1163/13210 train_time:276833ms step_avg:238.03ms |
| step:1164/13210 train_time:277061ms step_avg:238.02ms |
| step:1165/13210 train_time:277293ms step_avg:238.02ms |
| step:1166/13210 train_time:277520ms step_avg:238.01ms |
| step:1167/13210 train_time:277751ms step_avg:238.00ms |
| step:1168/13210 train_time:277978ms step_avg:238.00ms |
| step:1169/13210 train_time:278208ms step_avg:237.99ms |
| step:1170/13210 train_time:278436ms step_avg:237.98ms |
| step:1171/13210 train_time:278665ms step_avg:237.97ms |
| step:1172/13210 train_time:278892ms step_avg:237.96ms |
| step:1173/13210 train_time:279122ms step_avg:237.96ms |
| step:1174/13210 train_time:279349ms step_avg:237.95ms |
| step:1175/13210 train_time:279579ms step_avg:237.94ms |
| step:1176/13210 train_time:279807ms step_avg:237.93ms |
| step:1177/13210 train_time:280038ms step_avg:237.92ms |
| step:1178/13210 train_time:280265ms step_avg:237.92ms |
| step:1179/13210 train_time:280494ms step_avg:237.91ms |
| step:1180/13210 train_time:280722ms step_avg:237.90ms |
| step:1181/13210 train_time:280952ms step_avg:237.89ms |
| step:1182/13210 train_time:281179ms step_avg:237.88ms |
| step:1183/13210 train_time:281407ms step_avg:237.88ms |
| step:1184/13210 train_time:281636ms step_avg:237.87ms |
| step:1185/13210 train_time:281863ms step_avg:237.86ms |
| step:1186/13210 train_time:282093ms step_avg:237.85ms |
| step:1187/13210 train_time:282319ms step_avg:237.84ms |
| step:1188/13210 train_time:282549ms step_avg:237.84ms |
| step:1189/13210 train_time:282778ms step_avg:237.83ms |
| step:1190/13210 train_time:283004ms step_avg:237.82ms |
| step:1191/13210 train_time:283234ms step_avg:237.81ms |
| step:1192/13210 train_time:283462ms step_avg:237.80ms |
| step:1193/13210 train_time:283690ms step_avg:237.80ms |
| step:1194/13210 train_time:283920ms step_avg:237.79ms |
| step:1195/13210 train_time:284149ms step_avg:237.78ms |
| step:1196/13210 train_time:284379ms step_avg:237.78ms |
| step:1197/13210 train_time:284607ms step_avg:237.77ms |
| step:1198/13210 train_time:284836ms step_avg:237.76ms |
| step:1199/13210 train_time:285065ms step_avg:237.75ms |
| step:1200/13210 train_time:285292ms step_avg:237.74ms |
| step:1201/13210 train_time:285522ms step_avg:237.74ms |
| step:1202/13210 train_time:285748ms step_avg:237.73ms |
| step:1203/13210 train_time:285977ms step_avg:237.72ms |
| step:1204/13210 train_time:286205ms step_avg:237.71ms |
| step:1205/13210 train_time:286433ms step_avg:237.70ms |
| step:1206/13210 train_time:286660ms step_avg:237.70ms |
| step:1207/13210 train_time:286887ms step_avg:237.69ms |
| step:1208/13210 train_time:287115ms step_avg:237.68ms |
| step:1209/13210 train_time:287344ms step_avg:237.67ms |
| step:1210/13210 train_time:287570ms step_avg:237.66ms |
| step:1211/13210 train_time:287798ms step_avg:237.65ms |
| step:1212/13210 train_time:288027ms step_avg:237.65ms |
| step:1213/13210 train_time:288255ms step_avg:237.64ms |
| step:1214/13210 train_time:288486ms step_avg:237.63ms |
| step:1215/13210 train_time:288714ms step_avg:237.62ms |
| step:1216/13210 train_time:288942ms step_avg:237.62ms |
| step:1217/13210 train_time:289171ms step_avg:237.61ms |
| step:1218/13210 train_time:289399ms step_avg:237.60ms |
| step:1219/13210 train_time:289626ms step_avg:237.59ms |
| step:1220/13210 train_time:289854ms step_avg:237.59ms |
| step:1221/13210 train_time:290082ms step_avg:237.58ms |
| step:1222/13210 train_time:290310ms step_avg:237.57ms |
| step:1223/13210 train_time:290541ms step_avg:237.56ms |
| step:1224/13210 train_time:290769ms step_avg:237.56ms |
| step:1225/13210 train_time:290998ms step_avg:237.55ms |
| step:1226/13210 train_time:291225ms step_avg:237.54ms |
| step:1227/13210 train_time:291456ms step_avg:237.54ms |
| step:1228/13210 train_time:291685ms step_avg:237.53ms |
| step:1229/13210 train_time:291914ms step_avg:237.52ms |
| step:1230/13210 train_time:292143ms step_avg:237.51ms |
| step:1231/13210 train_time:292374ms step_avg:237.51ms |
| step:1232/13210 train_time:292606ms step_avg:237.50ms |
| step:1233/13210 train_time:292833ms step_avg:237.50ms |
| step:1234/13210 train_time:293063ms step_avg:237.49ms |
| step:1235/13210 train_time:293291ms step_avg:237.48ms |
| step:1236/13210 train_time:293519ms step_avg:237.47ms |
| step:1237/13210 train_time:293746ms step_avg:237.47ms |
| step:1238/13210 train_time:293972ms step_avg:237.46ms |
| step:1239/13210 train_time:294202ms step_avg:237.45ms |
| step:1240/13210 train_time:294428ms step_avg:237.44ms |
| step:1241/13210 train_time:294658ms step_avg:237.44ms |
| step:1242/13210 train_time:294887ms step_avg:237.43ms |
| step:1243/13210 train_time:295115ms step_avg:237.42ms |
| step:1244/13210 train_time:295345ms step_avg:237.42ms |
| step:1245/13210 train_time:295574ms step_avg:237.41ms |
| step:1246/13210 train_time:295802ms step_avg:237.40ms |
| step:1247/13210 train_time:296031ms step_avg:237.39ms |
| step:1248/13210 train_time:296259ms step_avg:237.39ms |
| step:1249/13210 train_time:296489ms step_avg:237.38ms |
| step:1250/13210 train_time:296717ms step_avg:237.37ms |
| step:1250/13210 val_loss:10.2671 train_time:296753ms step_avg:237.40ms |
| step:1251/13210 train_time:296949ms step_avg:237.37ms |
| step:1252/13210 train_time:297176ms step_avg:237.36ms |
| step:1253/13210 train_time:297405ms step_avg:237.35ms |
| step:1254/13210 train_time:297632ms step_avg:237.35ms |
| step:1255/13210 train_time:297860ms step_avg:237.34ms |
| step:1256/13210 train_time:298088ms step_avg:237.33ms |
| step:1257/13210 train_time:298316ms step_avg:237.32ms |
| step:1258/13210 train_time:298543ms step_avg:237.32ms |
| step:1259/13210 train_time:298774ms step_avg:237.31ms |
| step:1260/13210 train_time:299001ms step_avg:237.30ms |
| step:1261/13210 train_time:299231ms step_avg:237.30ms |
| step:1262/13210 train_time:299460ms step_avg:237.29ms |
| step:1263/13210 train_time:299687ms step_avg:237.28ms |
| step:1264/13210 train_time:299915ms step_avg:237.27ms |
| step:1265/13210 train_time:300143ms step_avg:237.27ms |
| step:1266/13210 train_time:300373ms step_avg:237.26ms |
| step:1267/13210 train_time:300602ms step_avg:237.26ms |
| step:1268/13210 train_time:300831ms step_avg:237.25ms |
| step:1269/13210 train_time:301060ms step_avg:237.24ms |
| step:1270/13210 train_time:301287ms step_avg:237.23ms |
| step:1271/13210 train_time:301517ms step_avg:237.23ms |
| step:1272/13210 train_time:301745ms step_avg:237.22ms |
| step:1273/13210 train_time:301976ms step_avg:237.22ms |
| step:1274/13210 train_time:302203ms step_avg:237.21ms |
| step:1275/13210 train_time:302434ms step_avg:237.20ms |
| step:1276/13210 train_time:302661ms step_avg:237.20ms |
| step:1277/13210 train_time:302889ms step_avg:237.19ms |
| step:1278/13210 train_time:303115ms step_avg:237.18ms |
| step:1279/13210 train_time:303344ms step_avg:237.17ms |
| step:1280/13210 train_time:303571ms step_avg:237.17ms |
| step:1281/13210 train_time:303801ms step_avg:237.16ms |
| step:1282/13210 train_time:304027ms step_avg:237.15ms |
| step:1283/13210 train_time:304254ms step_avg:237.14ms |
| step:1284/13210 train_time:304481ms step_avg:237.13ms |
| step:1285/13210 train_time:304710ms step_avg:237.13ms |
| step:1286/13210 train_time:304937ms step_avg:237.12ms |
| step:1287/13210 train_time:305165ms step_avg:237.11ms |
| step:1288/13210 train_time:305392ms step_avg:237.11ms |
| step:1289/13210 train_time:305621ms step_avg:237.10ms |
| step:1290/13210 train_time:305849ms step_avg:237.09ms |
| step:1291/13210 train_time:306077ms step_avg:237.08ms |
| step:1292/13210 train_time:306302ms step_avg:237.08ms |
| step:1293/13210 train_time:306530ms step_avg:237.07ms |
| step:1294/13210 train_time:306759ms step_avg:237.06ms |
| step:1295/13210 train_time:306987ms step_avg:237.06ms |
| step:1296/13210 train_time:307214ms step_avg:237.05ms |
| step:1297/13210 train_time:307441ms step_avg:237.04ms |
| step:1298/13210 train_time:307668ms step_avg:237.03ms |
| step:1299/13210 train_time:307896ms step_avg:237.03ms |
| step:1300/13210 train_time:308123ms step_avg:237.02ms |
| step:1301/13210 train_time:308350ms step_avg:237.01ms |
| step:1302/13210 train_time:308579ms step_avg:237.00ms |
| step:1303/13210 train_time:308808ms step_avg:237.00ms |
| step:1304/13210 train_time:309035ms step_avg:236.99ms |
| step:1305/13210 train_time:309262ms step_avg:236.98ms |
| step:1306/13210 train_time:309490ms step_avg:236.98ms |
| step:1307/13210 train_time:309718ms step_avg:236.97ms |
| step:1308/13210 train_time:309946ms step_avg:236.96ms |
| step:1309/13210 train_time:310172ms step_avg:236.95ms |
| step:1310/13210 train_time:310400ms step_avg:236.95ms |
| step:1311/13210 train_time:310629ms step_avg:236.94ms |
| step:1312/13210 train_time:310857ms step_avg:236.93ms |
| step:1313/13210 train_time:311085ms step_avg:236.93ms |
| step:1314/13210 train_time:311311ms step_avg:236.92ms |
| step:1315/13210 train_time:311539ms step_avg:236.91ms |
| step:1316/13210 train_time:311767ms step_avg:236.90ms |
| step:1317/13210 train_time:311994ms step_avg:236.90ms |
| step:1318/13210 train_time:312220ms step_avg:236.89ms |
| step:1319/13210 train_time:312448ms step_avg:236.88ms |
| step:1320/13210 train_time:312676ms step_avg:236.88ms |
| step:1321/13210 train_time:312904ms step_avg:236.87ms |
| step:1322/13210 train_time:313132ms step_avg:236.86ms |
| step:1323/13210 train_time:313359ms step_avg:236.85ms |
| step:1324/13210 train_time:313588ms step_avg:236.85ms |
| step:1325/13210 train_time:313818ms step_avg:236.84ms |
| step:1326/13210 train_time:314045ms step_avg:236.84ms |
| step:1327/13210 train_time:314272ms step_avg:236.83ms |
| step:1328/13210 train_time:314498ms step_avg:236.82ms |
| step:1329/13210 train_time:314726ms step_avg:236.81ms |
| step:1330/13210 train_time:314953ms step_avg:236.81ms |
| step:1331/13210 train_time:315182ms step_avg:236.80ms |
| step:1332/13210 train_time:315408ms step_avg:236.79ms |
| step:1333/13210 train_time:315635ms step_avg:236.79ms |
| step:1334/13210 train_time:315863ms step_avg:236.78ms |
| step:1335/13210 train_time:316092ms step_avg:236.77ms |
| step:1336/13210 train_time:316319ms step_avg:236.77ms |
| step:1337/13210 train_time:316546ms step_avg:236.76ms |
| step:1338/13210 train_time:316772ms step_avg:236.75ms |
| step:1339/13210 train_time:317002ms step_avg:236.75ms |
| step:1340/13210 train_time:317229ms step_avg:236.74ms |
| step:1341/13210 train_time:317457ms step_avg:236.73ms |
| step:1342/13210 train_time:317685ms step_avg:236.72ms |
| step:1343/13210 train_time:317914ms step_avg:236.72ms |
| step:1344/13210 train_time:318140ms step_avg:236.71ms |
| step:1345/13210 train_time:318369ms step_avg:236.71ms |
| step:1346/13210 train_time:318596ms step_avg:236.70ms |
| step:1347/13210 train_time:318824ms step_avg:236.69ms |
| step:1348/13210 train_time:319052ms step_avg:236.69ms |
| step:1349/13210 train_time:319280ms step_avg:236.68ms |
| step:1350/13210 train_time:319506ms step_avg:236.67ms |
| step:1351/13210 train_time:319735ms step_avg:236.67ms |
| step:1352/13210 train_time:319962ms step_avg:236.66ms |
| step:1353/13210 train_time:320190ms step_avg:236.65ms |
| step:1354/13210 train_time:320417ms step_avg:236.64ms |
| step:1355/13210 train_time:320645ms step_avg:236.64ms |
| step:1356/13210 train_time:320871ms step_avg:236.63ms |
| step:1357/13210 train_time:321100ms step_avg:236.63ms |
| step:1358/13210 train_time:321327ms step_avg:236.62ms |
| step:1359/13210 train_time:321555ms step_avg:236.61ms |
| step:1360/13210 train_time:321783ms step_avg:236.60ms |
| step:1361/13210 train_time:322010ms step_avg:236.60ms |
| step:1362/13210 train_time:322238ms step_avg:236.59ms |
| step:1363/13210 train_time:322465ms step_avg:236.58ms |
| step:1364/13210 train_time:322693ms step_avg:236.58ms |
| step:1365/13210 train_time:322920ms step_avg:236.57ms |
| step:1366/13210 train_time:323147ms step_avg:236.56ms |
| step:1367/13210 train_time:323376ms step_avg:236.56ms |
| step:1368/13210 train_time:323604ms step_avg:236.55ms |
| step:1369/13210 train_time:323832ms step_avg:236.55ms |
| step:1370/13210 train_time:324058ms step_avg:236.54ms |
| step:1371/13210 train_time:324286ms step_avg:236.53ms |
| step:1372/13210 train_time:324514ms step_avg:236.53ms |
| step:1373/13210 train_time:324743ms step_avg:236.52ms |
| step:1374/13210 train_time:324970ms step_avg:236.51ms |
| step:1375/13210 train_time:325196ms step_avg:236.51ms |
| step:1376/13210 train_time:325423ms step_avg:236.50ms |
| step:1377/13210 train_time:325651ms step_avg:236.49ms |
| step:1378/13210 train_time:325879ms step_avg:236.49ms |
| step:1379/13210 train_time:326106ms step_avg:236.48ms |
| step:1380/13210 train_time:326334ms step_avg:236.47ms |
| step:1381/13210 train_time:326562ms step_avg:236.47ms |
| step:1382/13210 train_time:326791ms step_avg:236.46ms |
| step:1383/13210 train_time:327019ms step_avg:236.46ms |
| step:1384/13210 train_time:327245ms step_avg:236.45ms |
| step:1385/13210 train_time:327472ms step_avg:236.44ms |
| step:1386/13210 train_time:327701ms step_avg:236.44ms |
| step:1387/13210 train_time:327929ms step_avg:236.43ms |
| step:1388/13210 train_time:328156ms step_avg:236.42ms |
| step:1389/13210 train_time:328384ms step_avg:236.42ms |
| step:1390/13210 train_time:328611ms step_avg:236.41ms |
| step:1391/13210 train_time:328840ms step_avg:236.41ms |
| step:1392/13210 train_time:329067ms step_avg:236.40ms |
| step:1393/13210 train_time:329294ms step_avg:236.39ms |
| step:1394/13210 train_time:329520ms step_avg:236.38ms |
| step:1395/13210 train_time:329748ms step_avg:236.38ms |
| step:1396/13210 train_time:329976ms step_avg:236.37ms |
| step:1397/13210 train_time:330203ms step_avg:236.37ms |
| step:1398/13210 train_time:330431ms step_avg:236.36ms |
| step:1399/13210 train_time:330658ms step_avg:236.35ms |
| step:1400/13210 train_time:330886ms step_avg:236.35ms |
| step:1401/13210 train_time:331115ms step_avg:236.34ms |
| step:1402/13210 train_time:331342ms step_avg:236.34ms |
| step:1403/13210 train_time:331572ms step_avg:236.33ms |
| step:1404/13210 train_time:331799ms step_avg:236.32ms |
| step:1405/13210 train_time:332027ms step_avg:236.32ms |
| step:1406/13210 train_time:332253ms step_avg:236.31ms |
| step:1407/13210 train_time:332480ms step_avg:236.30ms |
| step:1408/13210 train_time:332708ms step_avg:236.30ms |
| step:1409/13210 train_time:332935ms step_avg:236.29ms |
| step:1410/13210 train_time:333162ms step_avg:236.29ms |
| step:1411/13210 train_time:333390ms step_avg:236.28ms |
| step:1412/13210 train_time:333617ms step_avg:236.27ms |
| step:1413/13210 train_time:333846ms step_avg:236.27ms |
| step:1414/13210 train_time:334074ms step_avg:236.26ms |
| step:1415/13210 train_time:334301ms step_avg:236.26ms |
| step:1416/13210 train_time:334528ms step_avg:236.25ms |
| step:1417/13210 train_time:334757ms step_avg:236.24ms |
| step:1418/13210 train_time:334984ms step_avg:236.24ms |
| step:1419/13210 train_time:335212ms step_avg:236.23ms |
| step:1420/13210 train_time:335439ms step_avg:236.22ms |
| step:1421/13210 train_time:335666ms step_avg:236.22ms |
| step:1422/13210 train_time:335894ms step_avg:236.21ms |
| step:1423/13210 train_time:336122ms step_avg:236.21ms |
| step:1424/13210 train_time:336349ms step_avg:236.20ms |
| step:1425/13210 train_time:336576ms step_avg:236.19ms |
| step:1426/13210 train_time:336804ms step_avg:236.19ms |
| step:1427/13210 train_time:337032ms step_avg:236.18ms |
| step:1428/13210 train_time:337258ms step_avg:236.18ms |
| step:1429/13210 train_time:337486ms step_avg:236.17ms |
| step:1430/13210 train_time:337712ms step_avg:236.16ms |
| step:1431/13210 train_time:337940ms step_avg:236.16ms |
| step:1432/13210 train_time:338165ms step_avg:236.15ms |
| step:1433/13210 train_time:338393ms step_avg:236.14ms |
| step:1434/13210 train_time:338620ms step_avg:236.14ms |
| step:1435/13210 train_time:338848ms step_avg:236.13ms |
| step:1436/13210 train_time:339075ms step_avg:236.12ms |
| step:1437/13210 train_time:339303ms step_avg:236.12ms |
| step:1438/13210 train_time:339530ms step_avg:236.11ms |
| step:1439/13210 train_time:339760ms step_avg:236.11ms |
| step:1440/13210 train_time:339986ms step_avg:236.10ms |
| step:1441/13210 train_time:340213ms step_avg:236.09ms |
| step:1442/13210 train_time:340440ms step_avg:236.09ms |
| step:1443/13210 train_time:340667ms step_avg:236.08ms |
| step:1444/13210 train_time:340897ms step_avg:236.08ms |
| step:1445/13210 train_time:341124ms step_avg:236.07ms |
| step:1446/13210 train_time:341352ms step_avg:236.07ms |
| step:1447/13210 train_time:341580ms step_avg:236.06ms |
| step:1448/13210 train_time:341806ms step_avg:236.05ms |
| step:1449/13210 train_time:342033ms step_avg:236.05ms |
| step:1450/13210 train_time:342261ms step_avg:236.04ms |
| step:1451/13210 train_time:342488ms step_avg:236.04ms |
| step:1452/13210 train_time:342714ms step_avg:236.03ms |
| step:1453/13210 train_time:342942ms step_avg:236.02ms |
| step:1454/13210 train_time:343170ms step_avg:236.02ms |
| step:1455/13210 train_time:343398ms step_avg:236.01ms |
| step:1456/13210 train_time:343626ms step_avg:236.01ms |
| step:1457/13210 train_time:343853ms step_avg:236.00ms |
| step:1458/13210 train_time:344080ms step_avg:235.99ms |
| step:1459/13210 train_time:344309ms step_avg:235.99ms |
| step:1460/13210 train_time:344537ms step_avg:235.98ms |
| step:1461/13210 train_time:344764ms step_avg:235.98ms |
| step:1462/13210 train_time:344990ms step_avg:235.97ms |
| step:1463/13210 train_time:345219ms step_avg:235.97ms |
| step:1464/13210 train_time:345447ms step_avg:235.96ms |
| step:1465/13210 train_time:345676ms step_avg:235.96ms |
| step:1466/13210 train_time:345902ms step_avg:235.95ms |
| step:1467/13210 train_time:346129ms step_avg:235.94ms |
| step:1468/13210 train_time:346356ms step_avg:235.94ms |
| step:1469/13210 train_time:346584ms step_avg:235.93ms |
| step:1470/13210 train_time:346811ms step_avg:235.93ms |
| step:1471/13210 train_time:347040ms step_avg:235.92ms |
| step:1472/13210 train_time:347267ms step_avg:235.91ms |
| step:1473/13210 train_time:347495ms step_avg:235.91ms |
| step:1474/13210 train_time:347721ms step_avg:235.90ms |
| step:1475/13210 train_time:347948ms step_avg:235.90ms |
| step:1476/13210 train_time:348175ms step_avg:235.89ms |
| step:1477/13210 train_time:348403ms step_avg:235.89ms |
| step:1478/13210 train_time:348630ms step_avg:235.88ms |
| step:1479/13210 train_time:348859ms step_avg:235.87ms |
| step:1480/13210 train_time:349085ms step_avg:235.87ms |
| step:1481/13210 train_time:349314ms step_avg:235.86ms |
| step:1482/13210 train_time:349540ms step_avg:235.86ms |
| step:1483/13210 train_time:349768ms step_avg:235.85ms |
| step:1484/13210 train_time:349994ms step_avg:235.85ms |
| step:1485/13210 train_time:350223ms step_avg:235.84ms |
| step:1486/13210 train_time:350450ms step_avg:235.83ms |
| step:1487/13210 train_time:350679ms step_avg:235.83ms |
| step:1488/13210 train_time:350905ms step_avg:235.82ms |
| step:1489/13210 train_time:351133ms step_avg:235.82ms |
| step:1490/13210 train_time:351361ms step_avg:235.81ms |
| step:1491/13210 train_time:351589ms step_avg:235.81ms |
| step:1492/13210 train_time:351815ms step_avg:235.80ms |
| step:1493/13210 train_time:352044ms step_avg:235.80ms |
| step:1494/13210 train_time:352271ms step_avg:235.79ms |
| step:1495/13210 train_time:352499ms step_avg:235.79ms |
| step:1496/13210 train_time:352726ms step_avg:235.78ms |
| step:1497/13210 train_time:352953ms step_avg:235.77ms |
| step:1498/13210 train_time:353180ms step_avg:235.77ms |
| step:1499/13210 train_time:353407ms step_avg:235.76ms |
| step:1500/13210 train_time:353634ms step_avg:235.76ms |
| step:1500/13210 val_loss:10.2298 train_time:353670ms step_avg:235.78ms |
| step:1501/13210 train_time:357975ms step_avg:238.49ms |
| step:1502/13210 train_time:358199ms step_avg:238.48ms |
| step:1503/13210 train_time:358424ms step_avg:238.47ms |
| step:1504/13210 train_time:358649ms step_avg:238.46ms |
| step:1505/13210 train_time:358882ms step_avg:238.46ms |
| step:1506/13210 train_time:359107ms step_avg:238.45ms |
| step:1507/13210 train_time:359334ms step_avg:238.44ms |
| step:1508/13210 train_time:359559ms step_avg:238.43ms |
| step:1509/13210 train_time:359787ms step_avg:238.43ms |
| step:1510/13210 train_time:360014ms step_avg:238.42ms |
| step:1511/13210 train_time:360240ms step_avg:238.41ms |
| step:1512/13210 train_time:360467ms step_avg:238.40ms |
| step:1513/13210 train_time:360695ms step_avg:238.40ms |
| step:1514/13210 train_time:360923ms step_avg:238.39ms |
| step:1515/13210 train_time:361150ms step_avg:238.38ms |
| step:1516/13210 train_time:361375ms step_avg:238.37ms |
| step:1517/13210 train_time:361601ms step_avg:238.37ms |
| step:1518/13210 train_time:361829ms step_avg:238.36ms |
| step:1519/13210 train_time:362057ms step_avg:238.35ms |
| step:1520/13210 train_time:362284ms step_avg:238.34ms |
| step:1521/13210 train_time:362511ms step_avg:238.34ms |
| step:1522/13210 train_time:362736ms step_avg:238.33ms |
| step:1523/13210 train_time:362965ms step_avg:238.32ms |
| step:1524/13210 train_time:363193ms step_avg:238.32ms |
| step:1525/13210 train_time:363422ms step_avg:238.31ms |
| step:1526/13210 train_time:363649ms step_avg:238.30ms |
| step:1527/13210 train_time:363876ms step_avg:238.29ms |
| step:1528/13210 train_time:364104ms step_avg:238.29ms |
| step:1529/13210 train_time:364332ms step_avg:238.28ms |
| step:1530/13210 train_time:364559ms step_avg:238.27ms |
| step:1531/13210 train_time:364788ms step_avg:238.27ms |
| step:1532/13210 train_time:365015ms step_avg:238.26ms |
| step:1533/13210 train_time:365243ms step_avg:238.25ms |
| step:1534/13210 train_time:365468ms step_avg:238.25ms |
| step:1535/13210 train_time:365696ms step_avg:238.24ms |
| step:1536/13210 train_time:365924ms step_avg:238.23ms |
| step:1537/13210 train_time:366152ms step_avg:238.23ms |
| step:1538/13210 train_time:366379ms step_avg:238.22ms |
| step:1539/13210 train_time:366606ms step_avg:238.21ms |
| step:1540/13210 train_time:366833ms step_avg:238.20ms |
| step:1541/13210 train_time:367060ms step_avg:238.20ms |
| step:1542/13210 train_time:367287ms step_avg:238.19ms |
| step:1543/13210 train_time:367514ms step_avg:238.18ms |
| step:1544/13210 train_time:367739ms step_avg:238.17ms |
| step:1545/13210 train_time:367967ms step_avg:238.17ms |
| step:1546/13210 train_time:368194ms step_avg:238.16ms |
| step:1547/13210 train_time:368424ms step_avg:238.15ms |
| step:1548/13210 train_time:368652ms step_avg:238.15ms |
| step:1549/13210 train_time:368880ms step_avg:238.14ms |
| step:1550/13210 train_time:369108ms step_avg:238.13ms |
| step:1551/13210 train_time:369336ms step_avg:238.13ms |
| step:1552/13210 train_time:369563ms step_avg:238.12ms |
| step:1553/13210 train_time:369792ms step_avg:238.11ms |
| step:1554/13210 train_time:370018ms step_avg:238.11ms |
| step:1555/13210 train_time:370246ms step_avg:238.10ms |
| step:1556/13210 train_time:370473ms step_avg:238.09ms |
| step:1557/13210 train_time:370699ms step_avg:238.09ms |
| step:1558/13210 train_time:370929ms step_avg:238.08ms |
| step:1559/13210 train_time:371157ms step_avg:238.07ms |
| step:1560/13210 train_time:371385ms step_avg:238.07ms |
| step:1561/13210 train_time:371612ms step_avg:238.06ms |
| step:1562/13210 train_time:371840ms step_avg:238.05ms |
| step:1563/13210 train_time:372067ms step_avg:238.05ms |
| step:1564/13210 train_time:372295ms step_avg:238.04ms |
| step:1565/13210 train_time:372524ms step_avg:238.03ms |
| step:1566/13210 train_time:372751ms step_avg:238.03ms |
| step:1567/13210 train_time:372981ms step_avg:238.02ms |
| step:1568/13210 train_time:373207ms step_avg:238.01ms |
| step:1569/13210 train_time:373435ms step_avg:238.01ms |
| step:1570/13210 train_time:373662ms step_avg:238.00ms |
| step:1571/13210 train_time:373892ms step_avg:238.00ms |
| step:1572/13210 train_time:374120ms step_avg:237.99ms |
| step:1573/13210 train_time:374348ms step_avg:237.98ms |
| step:1574/13210 train_time:374574ms step_avg:237.98ms |
| step:1575/13210 train_time:374802ms step_avg:237.97ms |
| step:1576/13210 train_time:375030ms step_avg:237.96ms |
| step:1577/13210 train_time:375258ms step_avg:237.96ms |
| step:1578/13210 train_time:375486ms step_avg:237.95ms |
| step:1579/13210 train_time:375712ms step_avg:237.94ms |
| step:1580/13210 train_time:375940ms step_avg:237.94ms |
| step:1581/13210 train_time:376169ms step_avg:237.93ms |
| step:1582/13210 train_time:376396ms step_avg:237.92ms |
| step:1583/13210 train_time:376624ms step_avg:237.92ms |
| step:1584/13210 train_time:376851ms step_avg:237.91ms |
| step:1585/13210 train_time:377080ms step_avg:237.91ms |
| step:1586/13210 train_time:377308ms step_avg:237.90ms |
| step:1587/13210 train_time:377536ms step_avg:237.89ms |
| step:1588/13210 train_time:377763ms step_avg:237.89ms |
| step:1589/13210 train_time:377992ms step_avg:237.88ms |
| step:1590/13210 train_time:378218ms step_avg:237.87ms |
| step:1591/13210 train_time:378449ms step_avg:237.87ms |
| step:1592/13210 train_time:378675ms step_avg:237.86ms |
| step:1593/13210 train_time:378905ms step_avg:237.86ms |
| step:1594/13210 train_time:379132ms step_avg:237.85ms |
| step:1595/13210 train_time:379361ms step_avg:237.84ms |
| step:1596/13210 train_time:379587ms step_avg:237.84ms |
| step:1597/13210 train_time:379817ms step_avg:237.83ms |
| step:1598/13210 train_time:380045ms step_avg:237.83ms |
| step:1599/13210 train_time:380272ms step_avg:237.82ms |
| step:1600/13210 train_time:380501ms step_avg:237.81ms |
| step:1601/13210 train_time:380732ms step_avg:237.81ms |
| step:1602/13210 train_time:380959ms step_avg:237.80ms |
| step:1603/13210 train_time:381187ms step_avg:237.80ms |
| step:1604/13210 train_time:381413ms step_avg:237.79ms |
| step:1605/13210 train_time:381641ms step_avg:237.78ms |
| step:1606/13210 train_time:381868ms step_avg:237.78ms |
| step:1607/13210 train_time:382096ms step_avg:237.77ms |
| step:1608/13210 train_time:382324ms step_avg:237.76ms |
| step:1609/13210 train_time:382553ms step_avg:237.76ms |
| step:1610/13210 train_time:382783ms step_avg:237.75ms |
| step:1611/13210 train_time:383011ms step_avg:237.75ms |
| step:1612/13210 train_time:383240ms step_avg:237.74ms |
| step:1613/13210 train_time:383468ms step_avg:237.74ms |
| step:1614/13210 train_time:383695ms step_avg:237.73ms |
| step:1615/13210 train_time:383922ms step_avg:237.72ms |
| step:1616/13210 train_time:384150ms step_avg:237.72ms |
| step:1617/13210 train_time:384379ms step_avg:237.71ms |
| step:1618/13210 train_time:384606ms step_avg:237.70ms |
| step:1619/13210 train_time:384837ms step_avg:237.70ms |
| step:1620/13210 train_time:385063ms step_avg:237.69ms |
| step:1621/13210 train_time:385293ms step_avg:237.69ms |
| step:1622/13210 train_time:385520ms step_avg:237.68ms |
| step:1623/13210 train_time:385748ms step_avg:237.68ms |
| step:1624/13210 train_time:385977ms step_avg:237.67ms |
| step:1625/13210 train_time:386205ms step_avg:237.66ms |
| step:1626/13210 train_time:386434ms step_avg:237.66ms |
| step:1627/13210 train_time:386661ms step_avg:237.65ms |
| step:1628/13210 train_time:386888ms step_avg:237.65ms |
| step:1629/13210 train_time:387116ms step_avg:237.64ms |
| step:1630/13210 train_time:387345ms step_avg:237.63ms |
| step:1631/13210 train_time:387572ms step_avg:237.63ms |
| step:1632/13210 train_time:387800ms step_avg:237.62ms |
| step:1633/13210 train_time:388029ms step_avg:237.62ms |
| step:1634/13210 train_time:388258ms step_avg:237.61ms |
| step:1635/13210 train_time:388486ms step_avg:237.61ms |
| step:1636/13210 train_time:388714ms step_avg:237.60ms |
| step:1637/13210 train_time:388944ms step_avg:237.60ms |
| step:1638/13210 train_time:389172ms step_avg:237.59ms |
| step:1639/13210 train_time:389402ms step_avg:237.59ms |
| step:1640/13210 train_time:389630ms step_avg:237.58ms |
| step:1641/13210 train_time:389857ms step_avg:237.57ms |
| step:1642/13210 train_time:390084ms step_avg:237.57ms |
| step:1643/13210 train_time:390313ms step_avg:237.56ms |
| step:1644/13210 train_time:390542ms step_avg:237.56ms |
| step:1645/13210 train_time:390770ms step_avg:237.55ms |
| step:1646/13210 train_time:390997ms step_avg:237.54ms |
| step:1647/13210 train_time:391226ms step_avg:237.54ms |
| step:1648/13210 train_time:391454ms step_avg:237.53ms |
| step:1649/13210 train_time:391683ms step_avg:237.53ms |
| step:1650/13210 train_time:391909ms step_avg:237.52ms |
| step:1651/13210 train_time:392138ms step_avg:237.52ms |
| step:1652/13210 train_time:392365ms step_avg:237.51ms |
| step:1653/13210 train_time:392594ms step_avg:237.50ms |
| step:1654/13210 train_time:392823ms step_avg:237.50ms |
| step:1655/13210 train_time:393051ms step_avg:237.49ms |
| step:1656/13210 train_time:393277ms step_avg:237.49ms |
| step:1657/13210 train_time:393506ms step_avg:237.48ms |
| step:1658/13210 train_time:393734ms step_avg:237.48ms |
| step:1659/13210 train_time:393961ms step_avg:237.47ms |
| step:1660/13210 train_time:394190ms step_avg:237.46ms |
| step:1661/13210 train_time:394418ms step_avg:237.46ms |
| step:1662/13210 train_time:394646ms step_avg:237.45ms |
| step:1663/13210 train_time:394875ms step_avg:237.45ms |
| step:1664/13210 train_time:395102ms step_avg:237.44ms |
| step:1665/13210 train_time:395331ms step_avg:237.44ms |
| step:1666/13210 train_time:395557ms step_avg:237.43ms |
| step:1667/13210 train_time:395785ms step_avg:237.42ms |
| step:1668/13210 train_time:396014ms step_avg:237.42ms |
| step:1669/13210 train_time:396242ms step_avg:237.41ms |
| step:1670/13210 train_time:396471ms step_avg:237.41ms |
| step:1671/13210 train_time:396699ms step_avg:237.40ms |
| step:1672/13210 train_time:396927ms step_avg:237.40ms |
| step:1673/13210 train_time:397156ms step_avg:237.39ms |
| step:1674/13210 train_time:397385ms step_avg:237.39ms |
| step:1675/13210 train_time:397612ms step_avg:237.38ms |
| step:1676/13210 train_time:397838ms step_avg:237.37ms |
| step:1677/13210 train_time:398068ms step_avg:237.37ms |
| step:1678/13210 train_time:398295ms step_avg:237.36ms |
| step:1679/13210 train_time:398525ms step_avg:237.36ms |
| step:1680/13210 train_time:398751ms step_avg:237.35ms |
| step:1681/13210 train_time:398980ms step_avg:237.35ms |
| step:1682/13210 train_time:399207ms step_avg:237.34ms |
| step:1683/13210 train_time:399436ms step_avg:237.34ms |
| step:1684/13210 train_time:399664ms step_avg:237.33ms |
| step:1685/13210 train_time:399892ms step_avg:237.32ms |
| step:1686/13210 train_time:400121ms step_avg:237.32ms |
| step:1687/13210 train_time:400350ms step_avg:237.31ms |
| step:1688/13210 train_time:400577ms step_avg:237.31ms |
| step:1689/13210 train_time:400804ms step_avg:237.30ms |
| step:1690/13210 train_time:401031ms step_avg:237.30ms |
| step:1691/13210 train_time:401261ms step_avg:237.29ms |
| step:1692/13210 train_time:401488ms step_avg:237.29ms |
| step:1693/13210 train_time:401717ms step_avg:237.28ms |
| step:1694/13210 train_time:401944ms step_avg:237.27ms |
| step:1695/13210 train_time:402171ms step_avg:237.27ms |
| step:1696/13210 train_time:402399ms step_avg:237.26ms |
| step:1697/13210 train_time:402627ms step_avg:237.26ms |
| step:1698/13210 train_time:402853ms step_avg:237.25ms |
| step:1699/13210 train_time:403081ms step_avg:237.25ms |
| step:1700/13210 train_time:403308ms step_avg:237.24ms |
| step:1701/13210 train_time:403536ms step_avg:237.23ms |
| step:1702/13210 train_time:403764ms step_avg:237.23ms |
| step:1703/13210 train_time:403992ms step_avg:237.22ms |
| step:1704/13210 train_time:404219ms step_avg:237.22ms |
| step:1705/13210 train_time:404448ms step_avg:237.21ms |
| step:1706/13210 train_time:404675ms step_avg:237.21ms |
| step:1707/13210 train_time:404903ms step_avg:237.20ms |
| step:1708/13210 train_time:405130ms step_avg:237.20ms |
| step:1709/13210 train_time:405358ms step_avg:237.19ms |
| step:1710/13210 train_time:405585ms step_avg:237.18ms |
| step:1711/13210 train_time:405812ms step_avg:237.18ms |
| step:1712/13210 train_time:406040ms step_avg:237.17ms |
| step:1713/13210 train_time:406269ms step_avg:237.17ms |
| step:1714/13210 train_time:406495ms step_avg:237.16ms |
| step:1715/13210 train_time:406723ms step_avg:237.16ms |
| step:1716/13210 train_time:406949ms step_avg:237.15ms |
| step:1717/13210 train_time:407176ms step_avg:237.14ms |
| step:1718/13210 train_time:407402ms step_avg:237.14ms |
| step:1719/13210 train_time:407632ms step_avg:237.13ms |
| step:1720/13210 train_time:407859ms step_avg:237.13ms |
| step:1721/13210 train_time:408086ms step_avg:237.12ms |
| step:1722/13210 train_time:408314ms step_avg:237.12ms |
| step:1723/13210 train_time:408542ms step_avg:237.11ms |
| step:1724/13210 train_time:408769ms step_avg:237.11ms |
| step:1725/13210 train_time:408998ms step_avg:237.10ms |
| step:1726/13210 train_time:409227ms step_avg:237.10ms |
| step:1727/13210 train_time:409455ms step_avg:237.09ms |
| step:1728/13210 train_time:409682ms step_avg:237.08ms |
| step:1729/13210 train_time:409910ms step_avg:237.08ms |
| step:1730/13210 train_time:410138ms step_avg:237.07ms |
| step:1731/13210 train_time:410366ms step_avg:237.07ms |
| step:1732/13210 train_time:410592ms step_avg:237.06ms |
| step:1733/13210 train_time:410820ms step_avg:237.06ms |
| step:1734/13210 train_time:411049ms step_avg:237.05ms |
| step:1735/13210 train_time:411293ms step_avg:237.06ms |
| step:1736/13210 train_time:411564ms step_avg:237.08ms |
| step:1737/13210 train_time:411838ms step_avg:237.10ms |
| step:1738/13210 train_time:412107ms step_avg:237.12ms |
| step:1739/13210 train_time:412376ms step_avg:237.13ms |
| step:1740/13210 train_time:412645ms step_avg:237.15ms |
| step:1741/13210 train_time:412915ms step_avg:237.17ms |
| step:1742/13210 train_time:413172ms step_avg:237.18ms |
| step:1743/13210 train_time:413417ms step_avg:237.19ms |
| step:1744/13210 train_time:413642ms step_avg:237.18ms |
| step:1745/13210 train_time:413868ms step_avg:237.17ms |
| step:1746/13210 train_time:414099ms step_avg:237.17ms |
| step:1747/13210 train_time:414327ms step_avg:237.16ms |
| step:1748/13210 train_time:414551ms step_avg:237.16ms |
| step:1749/13210 train_time:414777ms step_avg:237.15ms |
| step:1750/13210 train_time:415005ms step_avg:237.15ms |
| step:1750/13210 val_loss:10.2200 train_time:415042ms step_avg:237.17ms |
| step:1751/13210 train_time:415236ms step_avg:237.14ms |
| step:1752/13210 train_time:415462ms step_avg:237.14ms |
| step:1753/13210 train_time:415690ms step_avg:237.13ms |
| step:1754/13210 train_time:415924ms step_avg:237.13ms |
| step:1755/13210 train_time:416146ms step_avg:237.12ms |
| step:1756/13210 train_time:416373ms step_avg:237.11ms |
| step:1757/13210 train_time:416601ms step_avg:237.11ms |
| step:1758/13210 train_time:416829ms step_avg:237.10ms |
| step:1759/13210 train_time:417057ms step_avg:237.10ms |
| step:1760/13210 train_time:417285ms step_avg:237.09ms |
| step:1761/13210 train_time:417513ms step_avg:237.09ms |
| step:1762/13210 train_time:417739ms step_avg:237.08ms |
| step:1763/13210 train_time:417967ms step_avg:237.08ms |
| step:1764/13210 train_time:418193ms step_avg:237.07ms |
| step:1765/13210 train_time:418422ms step_avg:237.07ms |
| step:1766/13210 train_time:418649ms step_avg:237.06ms |
| step:1767/13210 train_time:418876ms step_avg:237.05ms |
| step:1768/13210 train_time:419103ms step_avg:237.05ms |
| step:1769/13210 train_time:419331ms step_avg:237.04ms |
| step:1770/13210 train_time:419558ms step_avg:237.04ms |
| step:1771/13210 train_time:419787ms step_avg:237.03ms |
| step:1772/13210 train_time:420013ms step_avg:237.03ms |
| step:1773/13210 train_time:420241ms step_avg:237.02ms |
| step:1774/13210 train_time:420468ms step_avg:237.02ms |
| step:1775/13210 train_time:420697ms step_avg:237.01ms |
| step:1776/13210 train_time:420924ms step_avg:237.01ms |
| step:1777/13210 train_time:421153ms step_avg:237.00ms |
| step:1778/13210 train_time:421379ms step_avg:237.00ms |
| step:1779/13210 train_time:421606ms step_avg:236.99ms |
| step:1780/13210 train_time:421834ms step_avg:236.99ms |
| step:1781/13210 train_time:422063ms step_avg:236.98ms |
| step:1782/13210 train_time:422289ms step_avg:236.98ms |
| step:1783/13210 train_time:422517ms step_avg:236.97ms |
| step:1784/13210 train_time:422744ms step_avg:236.96ms |
| step:1785/13210 train_time:422973ms step_avg:236.96ms |
| step:1786/13210 train_time:423199ms step_avg:236.95ms |
| step:1787/13210 train_time:423426ms step_avg:236.95ms |
| step:1788/13210 train_time:423653ms step_avg:236.94ms |
| step:1789/13210 train_time:423880ms step_avg:236.94ms |
| step:1790/13210 train_time:424108ms step_avg:236.93ms |
| step:1791/13210 train_time:424337ms step_avg:236.93ms |
| step:1792/13210 train_time:424563ms step_avg:236.92ms |
| step:1793/13210 train_time:424791ms step_avg:236.92ms |
| step:1794/13210 train_time:425017ms step_avg:236.91ms |
| step:1795/13210 train_time:425245ms step_avg:236.91ms |
| step:1796/13210 train_time:425471ms step_avg:236.90ms |
| step:1797/13210 train_time:425700ms step_avg:236.89ms |
| step:1798/13210 train_time:425929ms step_avg:236.89ms |
| step:1799/13210 train_time:426157ms step_avg:236.89ms |
| step:1800/13210 train_time:426386ms step_avg:236.88ms |
| step:1801/13210 train_time:426613ms step_avg:236.88ms |
| step:1802/13210 train_time:426839ms step_avg:236.87ms |
| step:1803/13210 train_time:427066ms step_avg:236.86ms |
| step:1804/13210 train_time:427295ms step_avg:236.86ms |
| step:1805/13210 train_time:427525ms step_avg:236.86ms |
| step:1806/13210 train_time:427754ms step_avg:236.85ms |
| step:1807/13210 train_time:427982ms step_avg:236.85ms |
| step:1808/13210 train_time:428208ms step_avg:236.84ms |
| step:1809/13210 train_time:428437ms step_avg:236.84ms |
| step:1810/13210 train_time:428663ms step_avg:236.83ms |
| step:1811/13210 train_time:428893ms step_avg:236.83ms |
| step:1812/13210 train_time:429118ms step_avg:236.82ms |
| step:1813/13210 train_time:429345ms step_avg:236.81ms |
| step:1814/13210 train_time:429573ms step_avg:236.81ms |
| step:1815/13210 train_time:429801ms step_avg:236.81ms |
| step:1816/13210 train_time:430029ms step_avg:236.80ms |
| step:1817/13210 train_time:430256ms step_avg:236.79ms |
| step:1818/13210 train_time:430483ms step_avg:236.79ms |
| step:1819/13210 train_time:430712ms step_avg:236.79ms |
| step:1820/13210 train_time:430940ms step_avg:236.78ms |
| step:1821/13210 train_time:431167ms step_avg:236.77ms |
| step:1822/13210 train_time:431392ms step_avg:236.77ms |
| step:1823/13210 train_time:431618ms step_avg:236.76ms |
| step:1824/13210 train_time:431845ms step_avg:236.76ms |
| step:1825/13210 train_time:432074ms step_avg:236.75ms |
| step:1826/13210 train_time:432301ms step_avg:236.75ms |
| step:1827/13210 train_time:432529ms step_avg:236.74ms |
| step:1828/13210 train_time:432757ms step_avg:236.74ms |
| step:1829/13210 train_time:432983ms step_avg:236.73ms |
| step:1830/13210 train_time:433209ms step_avg:236.73ms |
| step:1831/13210 train_time:433436ms step_avg:236.72ms |
| step:1832/13210 train_time:433662ms step_avg:236.72ms |
| step:1833/13210 train_time:433892ms step_avg:236.71ms |
| step:1834/13210 train_time:434117ms step_avg:236.71ms |
| step:1835/13210 train_time:434359ms step_avg:236.71ms |
| step:1836/13210 train_time:434705ms step_avg:236.77ms |
| step:1837/13210 train_time:435110ms step_avg:236.86ms |
| step:1838/13210 train_time:435350ms step_avg:236.86ms |
| step:1839/13210 train_time:435576ms step_avg:236.85ms |
| step:1840/13210 train_time:435801ms step_avg:236.85ms |
| step:1841/13210 train_time:436036ms step_avg:236.85ms |
| step:1842/13210 train_time:436263ms step_avg:236.84ms |
| step:1843/13210 train_time:436490ms step_avg:236.84ms |
| step:1844/13210 train_time:436714ms step_avg:236.83ms |
| step:1845/13210 train_time:436945ms step_avg:236.83ms |
| step:1846/13210 train_time:437174ms step_avg:236.82ms |
| step:1847/13210 train_time:437400ms step_avg:236.82ms |
| step:1848/13210 train_time:437626ms step_avg:236.81ms |
| step:1849/13210 train_time:437853ms step_avg:236.81ms |
| step:1850/13210 train_time:438082ms step_avg:236.80ms |
| step:1851/13210 train_time:438311ms step_avg:236.80ms |
| step:1852/13210 train_time:438536ms step_avg:236.79ms |
| step:1853/13210 train_time:438764ms step_avg:236.79ms |
| step:1854/13210 train_time:438992ms step_avg:236.78ms |
| step:1855/13210 train_time:439220ms step_avg:236.78ms |
| step:1856/13210 train_time:439447ms step_avg:236.77ms |
| step:1857/13210 train_time:439674ms step_avg:236.77ms |
| step:1858/13210 train_time:439901ms step_avg:236.76ms |
| step:1859/13210 train_time:440130ms step_avg:236.76ms |
| step:1860/13210 train_time:440527ms step_avg:236.84ms |
| step:1861/13210 train_time:440752ms step_avg:236.84ms |
| step:1862/13210 train_time:440976ms step_avg:236.83ms |
| step:1863/13210 train_time:441202ms step_avg:236.82ms |
| step:1864/13210 train_time:441449ms step_avg:236.83ms |
| step:1865/13210 train_time:441721ms step_avg:236.85ms |
| step:1866/13210 train_time:441945ms step_avg:236.84ms |
| step:1867/13210 train_time:442171ms step_avg:236.84ms |
| step:1868/13210 train_time:442399ms step_avg:236.83ms |
| step:1869/13210 train_time:442632ms step_avg:236.83ms |
| step:1870/13210 train_time:442859ms step_avg:236.82ms |
| step:1871/13210 train_time:443085ms step_avg:236.82ms |
| step:1872/13210 train_time:443310ms step_avg:236.81ms |
| step:1873/13210 train_time:443541ms step_avg:236.81ms |
| step:1874/13210 train_time:443769ms step_avg:236.80ms |
| step:1875/13210 train_time:443997ms step_avg:236.80ms |
| step:1876/13210 train_time:444221ms step_avg:236.79ms |
| step:1877/13210 train_time:444449ms step_avg:236.79ms |
| step:1878/13210 train_time:444678ms step_avg:236.78ms |
| step:1879/13210 train_time:444906ms step_avg:236.78ms |
| step:1880/13210 train_time:445133ms step_avg:236.77ms |
| step:1881/13210 train_time:445360ms step_avg:236.77ms |
| step:1882/13210 train_time:445588ms step_avg:236.76ms |
| step:1883/13210 train_time:445816ms step_avg:236.76ms |
| step:1884/13210 train_time:446043ms step_avg:236.75ms |
| step:1885/13210 train_time:446270ms step_avg:236.75ms |
| step:1886/13210 train_time:446498ms step_avg:236.74ms |
| step:1887/13210 train_time:446728ms step_avg:236.74ms |
| step:1888/13210 train_time:446954ms step_avg:236.73ms |
| step:1889/13210 train_time:447182ms step_avg:236.73ms |
| step:1890/13210 train_time:447409ms step_avg:236.72ms |
| step:1891/13210 train_time:447637ms step_avg:236.72ms |
| step:1892/13210 train_time:447865ms step_avg:236.72ms |
| step:1893/13210 train_time:448093ms step_avg:236.71ms |
| step:1894/13210 train_time:448319ms step_avg:236.71ms |
| step:1895/13210 train_time:448547ms step_avg:236.70ms |
| step:1896/13210 train_time:448774ms step_avg:236.70ms |
| step:1897/13210 train_time:449002ms step_avg:236.69ms |
| step:1898/13210 train_time:449229ms step_avg:236.69ms |
| step:1899/13210 train_time:449458ms step_avg:236.68ms |
| step:1900/13210 train_time:449685ms step_avg:236.68ms |
| step:1901/13210 train_time:449915ms step_avg:236.67ms |
| step:1902/13210 train_time:450141ms step_avg:236.67ms |
| step:1903/13210 train_time:450368ms step_avg:236.66ms |
| step:1904/13210 train_time:450596ms step_avg:236.66ms |
| step:1905/13210 train_time:450824ms step_avg:236.65ms |
| step:1906/13210 train_time:451051ms step_avg:236.65ms |
| step:1907/13210 train_time:451278ms step_avg:236.64ms |
| step:1908/13210 train_time:451505ms step_avg:236.64ms |
| step:1909/13210 train_time:451733ms step_avg:236.63ms |
| step:1910/13210 train_time:451960ms step_avg:236.63ms |
| step:1911/13210 train_time:452186ms step_avg:236.62ms |
| step:1912/13210 train_time:452413ms step_avg:236.62ms |
| step:1913/13210 train_time:452643ms step_avg:236.61ms |
| step:1914/13210 train_time:452869ms step_avg:236.61ms |
| step:1915/13210 train_time:453096ms step_avg:236.60ms |
| step:1916/13210 train_time:453323ms step_avg:236.60ms |
| step:1917/13210 train_time:453550ms step_avg:236.59ms |
| step:1918/13210 train_time:453776ms step_avg:236.59ms |
| step:1919/13210 train_time:454004ms step_avg:236.58ms |
| step:1920/13210 train_time:454231ms step_avg:236.58ms |
| step:1921/13210 train_time:454458ms step_avg:236.57ms |
| step:1922/13210 train_time:454687ms step_avg:236.57ms |
| step:1923/13210 train_time:454915ms step_avg:236.57ms |
| step:1924/13210 train_time:455142ms step_avg:236.56ms |
| step:1925/13210 train_time:455369ms step_avg:236.56ms |
| step:1926/13210 train_time:455596ms step_avg:236.55ms |
| step:1927/13210 train_time:455823ms step_avg:236.55ms |
| step:1928/13210 train_time:456050ms step_avg:236.54ms |
| step:1929/13210 train_time:456277ms step_avg:236.54ms |
| step:1930/13210 train_time:456506ms step_avg:236.53ms |
| step:1931/13210 train_time:456734ms step_avg:236.53ms |
| step:1932/13210 train_time:456961ms step_avg:236.52ms |
| step:1933/13210 train_time:457188ms step_avg:236.52ms |
| step:1934/13210 train_time:457415ms step_avg:236.51ms |
| step:1935/13210 train_time:457643ms step_avg:236.51ms |
| step:1936/13210 train_time:457869ms step_avg:236.50ms |
| step:1937/13210 train_time:458099ms step_avg:236.50ms |
| step:1938/13210 train_time:458326ms step_avg:236.49ms |
| step:1939/13210 train_time:458554ms step_avg:236.49ms |
| step:1940/13210 train_time:458782ms step_avg:236.49ms |
| step:1941/13210 train_time:459009ms step_avg:236.48ms |
| step:1942/13210 train_time:459236ms step_avg:236.48ms |
| step:1943/13210 train_time:459465ms step_avg:236.47ms |
| step:1944/13210 train_time:459693ms step_avg:236.47ms |
| step:1945/13210 train_time:459922ms step_avg:236.46ms |
| step:1946/13210 train_time:460148ms step_avg:236.46ms |
| step:1947/13210 train_time:460376ms step_avg:236.45ms |
| step:1948/13210 train_time:460603ms step_avg:236.45ms |
| step:1949/13210 train_time:460830ms step_avg:236.44ms |
| step:1950/13210 train_time:461056ms step_avg:236.44ms |
| step:1951/13210 train_time:461284ms step_avg:236.43ms |
| step:1952/13210 train_time:461512ms step_avg:236.43ms |
| step:1953/13210 train_time:461740ms step_avg:236.43ms |
| step:1954/13210 train_time:461967ms step_avg:236.42ms |
| step:1955/13210 train_time:462195ms step_avg:236.42ms |
| step:1956/13210 train_time:462421ms step_avg:236.41ms |
| step:1957/13210 train_time:462649ms step_avg:236.41ms |
| step:1958/13210 train_time:462876ms step_avg:236.40ms |
| step:1959/13210 train_time:463103ms step_avg:236.40ms |
| step:1960/13210 train_time:463331ms step_avg:236.39ms |
| step:1961/13210 train_time:463559ms step_avg:236.39ms |
| step:1962/13210 train_time:463787ms step_avg:236.39ms |
| step:1963/13210 train_time:464015ms step_avg:236.38ms |
| step:1964/13210 train_time:464242ms step_avg:236.38ms |
| step:1965/13210 train_time:464470ms step_avg:236.37ms |
| step:1966/13210 train_time:464699ms step_avg:236.37ms |
| step:1967/13210 train_time:464927ms step_avg:236.36ms |
| step:1968/13210 train_time:465154ms step_avg:236.36ms |
| step:1969/13210 train_time:465383ms step_avg:236.35ms |
| step:1970/13210 train_time:465611ms step_avg:236.35ms |
| step:1971/13210 train_time:465840ms step_avg:236.35ms |
| step:1972/13210 train_time:466067ms step_avg:236.34ms |
| step:1973/13210 train_time:466296ms step_avg:236.34ms |
| step:1974/13210 train_time:466523ms step_avg:236.33ms |
| step:1975/13210 train_time:466752ms step_avg:236.33ms |
| step:1976/13210 train_time:466977ms step_avg:236.32ms |
| step:1977/13210 train_time:467204ms step_avg:236.32ms |
| step:1978/13210 train_time:467430ms step_avg:236.31ms |
| step:1979/13210 train_time:467657ms step_avg:236.31ms |
| step:1980/13210 train_time:467885ms step_avg:236.31ms |
| step:1981/13210 train_time:468113ms step_avg:236.30ms |
| step:1982/13210 train_time:468339ms step_avg:236.30ms |
| step:1983/13210 train_time:468567ms step_avg:236.29ms |
| step:1984/13210 train_time:468795ms step_avg:236.29ms |
| step:1985/13210 train_time:469023ms step_avg:236.28ms |
| step:1986/13210 train_time:469250ms step_avg:236.28ms |
| step:1987/13210 train_time:469477ms step_avg:236.27ms |
| step:1988/13210 train_time:469704ms step_avg:236.27ms |
| step:1989/13210 train_time:469932ms step_avg:236.27ms |
| step:1990/13210 train_time:470158ms step_avg:236.26ms |
| step:1991/13210 train_time:470387ms step_avg:236.26ms |
| step:1992/13210 train_time:470614ms step_avg:236.25ms |
| step:1993/13210 train_time:470843ms step_avg:236.25ms |
| step:1994/13210 train_time:471070ms step_avg:236.24ms |
| step:1995/13210 train_time:471298ms step_avg:236.24ms |
| step:1996/13210 train_time:471526ms step_avg:236.24ms |
|
|