add run results and update software versions

update red mlperf bert readme (#6969 )
qcom bench bind to perf cluster (#6996 )
2026-08-24 22:46:07 +00:00 · 2024-10-11 10:32:29 -04:00 · 2024-10-11 13:08:06 +03:00 · 2024-10-11 12:21:52 +03:00 · 2024-10-11 11:33:16 +03:00 · 2024-10-11 11:13:46 +03:00
279 changed files with 111299 additions and 43244 deletions
@@ -1,7 +1,6 @@
 name: Benchmarks
 env:
  # TODO: this rescheduling makes gpt2, mixtral and llama unjitted slower
-  COMPARE_SCHEDULE: "0"
  RUN_PROCESS_REPLAY: "1"
  ASSERT_PROCESS_REPLAY: "0"
  PYTHONPATH: .
@@ -45,6 +44,8 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
    - name: Run Stable Diffusion
      run: JIT=2 python3 examples/stable_diffusion.py --seed 0 --noshow --timing | tee sd.txt
    - name: Run Stable Diffusion with fp16
@@ -148,6 +149,8 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
    - name: Run model inference benchmark
      run: NV=1 NOCLANG=1 python3 test/external/external_model_benchmark.py
    - name: Test speed vs torch
@@ -166,10 +169,6 @@ jobs:
      run: NV=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py | tee matmul_nv.txt
    - name: Run Tensor Core GEMM (NV) with BEAM
      run: BEAM=4 NV=1 HALF=1 IGNORE_BEAM_CACHE=1 DEBUG=2 python3 extra/gemm/simple_matmul.py
-    - name: Fuzz Padded Tensor Core GEMM (NV)
-      run: NV=1 M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
-    - name: Fuzz Padded Tensor Core GEMM (PTX)
-      run: NV=1 PTX=1 M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
    - name: Run Stable Diffusion
      run: NV=1 python3 examples/stable_diffusion.py --seed 0 --noshow --timing | tee sd.txt
    - name: Run SDXL
@@ -256,6 +255,12 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
+    - name: Fuzz Padded Tensor Core GEMM (NV)
+      run: NV=1 M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
+    - name: Fuzz Padded Tensor Core GEMM (PTX)
+      run: NV=1 PTX=1 M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py
    - name: Train MNIST
      run: time PYTHONPATH=. NV=1 TARGET_EVAL_ACC_PCT=97.3 python3 examples/beautiful_mnist.py | tee beautiful_mnist.txt
    - name: Run 10 CIFAR training steps
@@ -318,6 +323,8 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
    - name: Show off tinybox
      run: /opt/rocm/bin/rocm-bandwidth-test
    # TODO: unstable on AMD
@@ -420,6 +427,8 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
    - name: Train MNIST
      run: time PYTHONPATH=. AMD=1 TARGET_EVAL_ACC_PCT=97.3 python3 examples/beautiful_mnist.py | tee beautiful_mnist.txt
    - name: Run 10 CIFAR training steps
@@ -471,20 +480,28 @@ jobs:
      run: |
        echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV
        rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal
+    - name: reset process replay
+      run: test/external/process_replay/reset.py
    - name: openpilot compile 0.9.4
-      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 GPU=1 python examples/openpilot/compile2.py | tee openpilot_compile_0_9_4.txt
+      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python examples/openpilot/compile2.py | tee openpilot_compile_0_9_4.txt
    - name: openpilot compile 0.9.7
-      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 GPU=1 python examples/openpilot/compile2.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_compile_0_9_7.txt
+      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python examples/openpilot/compile2.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_compile_0_9_7.txt
    - name: validate openpilot 0.9.7
-      run: PYTHONPATH=. FLOAT16=0 IMAGE=2 GPU=1 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_7.txt
+      run: PYTHONPATH=. FLOAT16=0 IMAGE=2 QCOM=1 taskset -c 4-7 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_7.txt
    - name: benchmark openpilot 0.9.4
-      run: PYTHONPATH=. GPU=1 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.4/selfdrive/modeld/models/supercombo.onnx | tee openpilot_0_9_4.txt
+      run: PYTHONPATH=. QCOM=1 taskset -c 4-7 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.4/selfdrive/modeld/models/supercombo.onnx | tee openpilot_0_9_4.txt
    - name: benchmark openpilot 0.9.7
-      run: PYTHONPATH=. GPU=1 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_0_9_7.txt
+      run: PYTHONPATH=. QCOM=1 taskset -c 4-7 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_0_9_7.txt
    - name: benchmark openpilot w IMAGE=2 0.9.4
-      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 GPU=1 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.4/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_4.txt
+      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.4/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_4.txt
    - name: benchmark openpilot w IMAGE=2 0.9.7
-      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 GPU=1 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_7.txt
+      run: PYTHONPATH=. NOLOCALS=1 FLOAT16=1 IMAGE=2 QCOM=1 taskset -c 4-7 python3 test/external/external_benchmark_openpilot.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx | tee openpilot_image_0_9_7.txt
+    - name: openpilot compile3 0.9.7
+      run: PYTHONPATH="." QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx
+    - name: openpilot compile3 0.9.7+ tomb raider
+      run: PYTHONPATH="." QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/e8bea2c78ffa92685ece511e9b554122aaf1a79d/selfdrive/modeld/models/supercombo.onnx
+    - name: openpilot dmonitoring compile3 0.9.7
+      run: PYTHONPATH="." QCOM=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/dmonitoring_model.onnx
    - name: Run process replay tests
      run: cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
    - uses: actions/upload-artifact@v4
@@ -1,7 +1,7 @@
 name: Unit Tests
 env:
  # increment this when downloads substantially change to avoid the internet
-  DOWNLOAD_CACHE_VERSION: '5'
+  DOWNLOAD_CACHE_VERSION: '6'
  RUN_PROCESS_REPLAY: 1
  GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
  PYTHONPATH: .
@@ -38,6 +38,8 @@ jobs:
        IMAGE=2 PYTHON=1 python3 test/test_ops.py TestOps.test_simple_conv2d
    - name: Test emulated METAL tensor cores
      run: DEBUG=2 EMULATE_METAL=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_big_gemm
+    - name: Test emulated AMX tensor cores
+      run: PYTHONPATH=. DEBUG=2 AMX=1 EMULATE_AMX=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_gemm
    - name: Test emulated AMD tensor cores
      run: |
        PYTHONPATH=. DEBUG=2 EMULATE_AMD=1 FORWARD_ONLY=1 PYTHON=1 N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
@@ -46,16 +48,29 @@ jobs:
        PYTHONPATH=. DEBUG=2 EMULATE_AMD=1 FORWARD_ONLY=1 PYTHON=1 N=64 HALF=1 ACC_HALF=1 python3 ./extra/gemm/simple_matmul.py
    - name: Test emulated CUDA tensor cores
      run: DEBUG=2 EMULATE_CUDA=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_gemm_fp16
+    - name: Test emulated INTEL OpenCL tensor cores
+      run: DEBUG=2 EMULATE_INTEL=1 FORWARD_ONLY=1 PYTHON=1 HALF=1 N=64 python3 ./extra/gemm/simple_matmul.py
    - name: Full test tensor cores
      run: |
        PYTHONPATH=. DEBUG=2 EMULATE_METAL=1 FORWARD_ONLY=1 PYTHON=1 python3 ./test/test_linearizer.py TestLinearizer.test_tensor_cores
        PYTHONPATH=. DEBUG=2 EMULATE_AMD=1 FORWARD_ONLY=1 PYTHON=1 python3 ./test/test_linearizer.py TestLinearizer.test_tensor_cores
        PYTHONPATH=. DEBUG=2 EMULATE_CUDA=1 FORWARD_ONLY=1 PYTHON=1 python3 ./test/test_linearizer.py TestLinearizer.test_tensor_cores
+        PYTHONPATH=. DEBUG=2 EMULATE_INTEL=1 FORWARD_ONLY=1 PYTHON=1 python3 ./test/test_linearizer.py TestLinearizer.test_tensor_cores
+        PYTHONPATH=. DEBUG=2 AMX=1 EMULATE_AMX=1 FORWARD_ONLY=1 PYTHON=1 python3 ./test/test_linearizer.py TestLinearizer.test_tensor_cores
    - name: Test tensor cores (TC=3)
      run: |
        TC=3 DEBUG=3 EMULATE_METAL=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_gemm
-        TC=3 PYTHONPATH=. DEBUG=3 EMULATE_AMD=1 FORWARD_ONLY=1 PYTHON=1 N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
+        TC=3 PYTHONPATH=. DEBUG=3 EMULATE_AMD=1 PYTHON=1 N=16 HALF=1 ACC_HALF=0 python3 ./extra/gemm/simple_matmul.py
        TC=3 DEBUG=3 EMULATE_CUDA=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_gemm_fp16
+        TC=3 PYTHONPATH=. DEBUG=3 EMULATE_INTEL=1 PYTHON=1 N=16 HALF=1 python3 ./extra/gemm/simple_matmul.py
+        TC=3 PYTHONPATH=. DEBUG=3 AMX=1 EMULATE_AMX=1 FORWARD_ONLY=1 PYTHON=1 python3 test/test_ops.py TestOps.test_gemm
+    - name: Test device flop counts
+      run: |
+        PYTHONPATH=. DEBUG=2 EMULATE_METAL=1 PYTHON=1 python3 ./test/test_uops_stats.py TestUOpsStatsMatmulHalf
+        PYTHONPATH=. DEBUG=2 EMULATE_AMD=1 PYTHON=1 python3 ./test/test_uops_stats.py TestUOpsStatsMatmulHalf
+        PYTHONPATH=. DEBUG=2 EMULATE_CUDA=1 PYTHON=1 python3 ./test/test_uops_stats.py TestUOpsStatsMatmulHalf
+        PYTHONPATH=. DEBUG=2 EMULATE_INTEL=1 PYTHON=1 python3 ./test/test_uops_stats.py TestUOpsStatsMatmulHalf
+        PYTHONPATH=. DEBUG=2 AMX=1 EMULATE_AMX=1 PYTHON=1 python3 ./test/test_uops_stats.py TestUOpsStats.test_simple_matmul
    - name: Test dtype with Python emulator
      run: DEBUG=1 PYTHONPATH=. PYTHON=1 python3 -m pytest -n=auto test/test_dtype.py test/test_dtype_alu.py
    - name: Test ops with Python emulator
@@ -115,8 +130,6 @@ jobs:
      run: |
        PYTHONPATH="." python test/external/fuzz_shapetracker.py
        PYTHONPATH="." python test/external/fuzz_shapetracker_math.py
-    - name: Test to_movement_ops
-      run: PYTHONPATH="." python extra/to_movement_ops.py
    - name: Use as an external package
      run: |
        mkdir $HOME/test_external_dir
@@ -125,10 +138,21 @@ jobs:
        source venv/bin/activate
        pip install $GITHUB_WORKSPACE
        python -c "from tinygrad.tensor import Tensor; print(Tensor([1,2,3,4,5]))"
+        pip install mypy
+        mypy -c "from tinygrad.tensor import Tensor; print(Tensor([1,2,3,4,5]))"
+    - name: Run beautiful_mnist without numpy
+      run: |
+        mkdir $HOME/test_no_numpy_dir
+        cd $HOME/test_no_numpy_dir
+        python -m venv venv
+        source venv/bin/activate
+        pip install $GITHUB_WORKSPACE
+        cp $GITHUB_WORKSPACE/examples/beautiful_mnist.py .
+        PYTHONPATH=$GITHUB_WORKSPACE BS=2 STEPS=10 python beautiful_mnist.py
    - name: Test DEBUG
      run: DEBUG=100 python3 -c "from tinygrad import Tensor; N = 1024; a, b = Tensor.rand(N, N), Tensor.rand(N, N); c = (a.reshape(N, 1, N) * b.T.reshape(1, N, N)).sum(axis=2); print((c.numpy() - (a.numpy() @ b.numpy())).mean())"
-    - name: Repo line count <8800 lines
-      run: MAX_LINE_COUNT=8800 python sz.py
+    - name: Repo line count <= 9999 lines
+      run: MAX_LINE_COUNT=9999 python sz.py

  testopencl:
    strategy:
@@ -182,7 +206,7 @@ jobs:
      - if: ${{ matrix.task == 'optimage' }}
        name: Test openpilot model compile and size
        run: |
-          PYTHONPATH="." DEBUG=2 ALLOWED_KERNEL_COUNT=208 FLOAT16=1 DEBUGCL=1 GPU=1 IMAGE=2 python examples/openpilot/compile2.py
+          PYTHONPATH="." DEBUG=2 ALLOWED_KERNEL_COUNT=208 ALLOWED_GATED_READ_IMAGE=13 FLOAT16=1 DEBUGCL=1 GPU=1 IMAGE=2 python examples/openpilot/compile2.py
          python -c 'import os; assert os.path.getsize("/tmp/output.thneed") < 100_000_000'
      - if: ${{ matrix.task == 'optimage' }}
        name: Test openpilot model correctness (float32)
@@ -205,6 +229,9 @@ jobs:
      - if: ${{ matrix.task == 'onnx' }}
        name: Test ONNX (CLANG)
        run: CLANG=1 python -m pytest -n=auto test/external/external_test_onnx_backend.py --durations=20
+      - if: ${{ matrix.task == 'onnx' }}
+        name: Run CLOUD=1 Test
+        run: CLOUDDEV=CLANG CLOUD=1 python3 test/test_ops.py TestOps.test_tiny_add
      - if: ${{ matrix.task == 'onnx' }}
        name: Test Action Space
        run: PYTHONPATH="." GPU=1 python3 extra/optimization/get_action_space.py
@@ -223,9 +250,6 @@ jobs:
      - if: ${{ matrix.task == 'onnx' }}
        name: Test MLPerf datasets
        run: GPU=1 python -m pytest -n=auto test/external/external_test_datasets.py --durations=20
-      - if: ${{ matrix.task == 'onnx' }}
-        name: Test THREEFRY
-        run: PYTHONPATH=. THREEFRY=1 GPU=1 python3 -m pytest test/test_randomness.py test/test_jit.py --durations=20
      - if: ${{ matrix.task == 'onnx' }}
        name: Run handcode_opt
        run: PYTHONPATH=. MODEL=resnet GPU=1 DEBUG=1 BS=4 HALF=0 python3 examples/handcode_opt.py
@@ -326,12 +350,11 @@ jobs:
      run: FUZZ_SCHEDULE=1 FUZZ_SCHEDULE_MAX_PATHS=5 python -m pytest test/models/test_train.py test/models/test_end2end.py
    - name: Run TRANSCENDENTAL math
      run: TRANSCENDENTAL=2 python -m pytest -n=auto test/test_ops.py::TestOps::test_sin test/test_ops.py::TestOps::test_cos test/test_ops.py::TestOps::test_tan test/test_ops.py::TestOps::test_exp test/test_ops.py::TestOps::test_log --durations=20
-      # TODO: this is timing out
-      #- name: Run process replay tests
-      #run: |
-      #  export PR_TITLE=$(jq -r .pull_request.title "$GITHUB_EVENT_PATH")
-      #  export COMMIT_MESSAGE=$(git show -s --format=%B ${{ github.event.pull_request.head.sha }})
-      #  cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
+    - name: Run process replay tests
+      run: |
+        export PR_TITLE=$(jq -r .pull_request.title "$GITHUB_EVENT_PATH")
+        export COMMIT_MESSAGE=$(git show -s --format=%B ${{ github.event.pull_request.head.sha }})
+        cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py

 #  testwebgl:
 #    name: WebGL Tests
@@ -422,7 +445,7 @@ jobs:
          cache-name: cache-gpuocelot-build
        with:
          path: ${{ github.workspace }}/gpuocelot/ocelot
-          key: ubuntu22.04-gpuocelot-4524e34adb7eaccc6f71262f2e21d7052bb17c2f-rebuild-7
+          key: ubuntu22.04-gpuocelot-4524e34adb7eaccc6f71262f2e21d7052bb17c2f-rebuild-8
      - name: Clone/compile gpuocelot
        if: (matrix.backend == 'ptx' || matrix.backend == 'triton' || matrix.backend == 'nv') && steps.cache-build.outputs.cache-hit != 'true'
        run: |
@@ -518,9 +541,6 @@ jobs:
        run: TRANSCENDENTAL=2 python -m pytest -n=auto test/test_ops.py::TestOps::test_sin test/test_ops.py::TestOps::test_cos test/test_ops.py::TestOps::test_tan test/test_ops.py::TestOps::test_exp test/test_ops.py::TestOps::test_log --durations=20
      - name: Run process replay tests
        run: |
-          if [ "${{ matrix.backend }}" == "amd" ] && [ "${GITHUB_REF_NAME}" != "master" ]; then
-            MAX_DIFF_PCT=1 RUN_PROCESS_REPLAY=0 test/external/process_replay/test_process_replay.sh
-          fi
          export PR_TITLE=$(jq -r .pull_request.title "$GITHUB_EVENT_PATH")
          export COMMIT_MESSAGE=$(git show -s --format=%B ${{ github.event.pull_request.head.sha }})
          cp test/external/process_replay/process_replay.py ./process_replay.py && git fetch origin master && git -c advice.detachedHead=false checkout origin/master && PYTHONPATH=. python3 process_replay.py
@@ -27,13 +27,13 @@ repos:
        pass_filenames: false
      - id: devicetests
        name: select GPU tests
-        entry: env GPU=1 PYTHONPATH="." pytest test/test_uops.py test/test_custom_function.py test/test_search.py
+        entry: env GPU=1 PYTHONPATH="." pytest test/test_uops.py test/test_search.py
        language: system
        always_run: true
        pass_filenames: false
      - id: tests
        name: subset of tests
-        entry: env PYTHONPATH="." python3 -m pytest -n=4 test/unit/ test/test_ops.py test/test_dtype.py test/test_schedule.py test/test_custom_function.py test/test_assign.py test/test_symbolic_shapetracker.py
+        entry: env PYTHONPATH="." python3 -m pytest -n=4 test/unit/ test/test_ops.py test/test_dtype.py test/test_schedule.py test/test_assign.py test/test_symbolic_shapetracker.py
        language: system
        always_run: true
        pass_filenames: false
@@ -87,9 +87,12 @@ tinygrad already supports numerous accelerators, including:
 - [x] [CUDA](tinygrad/runtime/ops_cuda.py)
 - [x] [AMD](tinygrad/runtime/ops_amd.py)
 - [x] [NV](tinygrad/runtime/ops_nv.py)
+- [x] [QCOM](tinygrad/runtime/ops_qcom.py)

 And it is easy to add more! Your accelerator of choice only needs to support a total of ~25 low level ops.

+To check default accelerator run: `python3 -c "from tinygrad import Device; print(Device.DEFAULT)"`
+
 ## Installation

 The current recommended way to install tinygrad is from source.
@@ -175,4 +178,4 @@ python3 -m pytest test/                 # whole test suite

 #### Process replay tests

-[Process replay](https://github.com/tinygrad/tinygrad/blob/master/test/external/process_replay/process_replay.py) compares your PR's generated kernels against master. If your PR is a refactor or speedup without any expected behavior change, It should include [run_process_replay] in the PR title, [example](https://github.com/tinygrad/tinygrad/pull/4995). Note that you should keep your branch up-to-date with master.
+[Process replay](https://github.com/tinygrad/tinygrad/blob/master/test/external/process_replay/process_replay.py) compares your PR's generated kernels against master. If your PR is a refactor or speedup without any expected behavior change, It should include [pr] in the pull request title, [example](https://github.com/tinygrad/tinygrad/pull/4995). Note that you should keep your branch up-to-date with master.
@@ -7,7 +7,7 @@ if [[ ! $(clang2py -V) ]]; then
  sudo apt-get install -y --no-install-recommends clang
  pip install --upgrade pip setuptools
  pip install clang==14.0.6
-  git clone https://github.com/geohot/ctypeslib.git
+  git clone https://github.com/nimlgen/ctypeslib.git
  cd ctypeslib
  pip install --user .
  clang2py -V
@@ -39,56 +39,6 @@ def _try_dlopen_$name():
 EOF
 }

-process_cdefines() {
-  local input_file="$1"
-  local output_file="$2"
-
-  sed -E '
-    # Remove single-line comments
-    s/[[:space:]]*\/\*.*\*\///g
-
-    # Remove multi-line comments
-    /\/\*/,/\*\//d
-
-    /.*DT_MIPS_NUM.*/d
-
-    # Remove lines ending with backslash (multi-line macros)
-    /\\$/d
-
-    # Convert C integer literals (remove U suffix)
-    s/\b([0-9]+)U\b/\1/g
-
-    # Convert C types to Python ctypes
-    s/\bunsigned char\b/ctypes.c_ubyte/g
-    s/\bsigned char\b/ctypes.c_byte/g
-    s/\bunsigned short\b/ctypes.c_ushort/g
-    s/\bshort\b/ctypes.c_short/g
-    s/\bunsigned int\b/ctypes.c_uint/g
-    s/\bint\b/ctypes.c_int/g
-    s/\bunsigned long\b/ctypes.c_ulong/g
-    s/\blong\b/ctypes.c_long/g
-    s/\bfloat\b/ctypes.c_float/g
-    s/\bdouble\b/ctypes.c_double/g
-
-    # Function-like macros with parameters
-    /^#define[[:space:]]+([[:alnum:]_]+)[[:space:]]*\(([^)]*)\)[[:space:]]+(.+)/ {
-      s//def \1(\2): return \3/
-      p
-      d
-    }
-
-    # Simple #define statements (including those with parentheses)
-    /^#define[[:space:]]+([[:alnum:]_]+)[[:space:]]+(.+)/ {
-      s//\1 = \2/
-      p
-      d
-    }
-
-    # Drop all other lines
-    d
-  ' "$input_file" >> "$output_file"
-}
-
 generate_opencl() {
  clang2py /usr/include/CL/cl.h -o $BASE/opencl.py -l /usr/lib/x86_64-linux-gnu/libOpenCL.so.1 -k cdefstum
  fixup $BASE/opencl.py
@@ -126,11 +76,7 @@ generate_comgr() {

 generate_kfd() {
  clang2py /usr/include/linux/kfd_ioctl.h -o $BASE/kfd.py -k cdefstum
-  awk '/^#define AMDKFD_IOC_/ { if ($0 ~ /\\$/) { getline nextline; $0 = $0 nextline }
-    if (match($0, /AMDKFD_IOC_([A-Z_]+).*AMDKFD_(IOW?R?)\(0x([0-9A-F]+),.*struct ([a-z_]+)/, arr)) {
-      print "AMDKFD_IOC_" arr[1] " = (\"" arr[2] "\", 0x" arr[3] ", struct_" arr[4] ")"
-    }
-  }' /usr/include/linux/kfd_ioctl.h >> $BASE/kfd.py
+
  fixup $BASE/kfd.py
  sed -i "s\import ctypes\import ctypes, os\g" $BASE/kfd.py
  python3 -c "import tinygrad.runtime.autogen.kfd"
@@ -164,7 +110,7 @@ generate_nv() {
    popd
  fi

-  clang2py \
+  clang2py -k cdefstum \
    extra/nv_gpu_driver/clc6c0qmd.h \
    $NVKERN_SRC/src/common/sdk/nvidia/inc/class/cl0080.h \
    $NVKERN_SRC/src/common/sdk/nvidia/inc/class/cl2080_notification.h \
@@ -191,7 +137,7 @@ generate_nv() {
    $NVKERN_SRC/src/common/sdk/nvidia/inc/ctrl/ctrlcb33.h \
    $NVKERN_SRC/src/common/sdk/nvidia/inc/ctrl/ctrla06c.h \
    --clang-args="-include $NVKERN_SRC/src/common/sdk/nvidia/inc/nvtypes.h -I$NVKERN_SRC/src/common/inc -I$NVKERN_SRC/kernel-open/nvidia-uvm -I$NVKERN_SRC/kernel-open/common/inc -I$NVKERN_SRC/src/common/sdk/nvidia/inc -I$NVKERN_SRC/src/nvidia/arch/nvalloc/unix/include -I$NVKERN_SRC/src/common/sdk/nvidia/inc/ctrl" \
-    -o $BASE/nv_gpu.py -k cdefstum
+    -o $BASE/nv_gpu.py
  fixup $BASE/nv_gpu.py
  sed -i "s\(0000000001)\1\g" $BASE/nv_gpu.py
  sed -i "s\import ctypes\import ctypes, os\g" $BASE/nv_gpu.py
@@ -211,23 +157,15 @@ nv_status_codes = {}

 generate_amd() {
  # clang2py broken when pass -x c++ to prev headers
-  clang2py extra/hip_gpu_driver/sdma_registers.h \
+  clang2py -k cdefstum \
+    extra/hip_gpu_driver/sdma_registers.h \
+    extra/hip_gpu_driver/nvd.h \
+    extra/hip_gpu_driver/sdma_v6_0_0_pkt_open.h \
+    extra/hip_gpu_driver/gc_11_0_0_offset.h \
+    extra/hip_gpu_driver/gc_10_3_0_offset.h \
    --clang-args="-I/opt/rocm/include -x c++" \
    -o $BASE/amd_gpu.py

-  sed 's/^\(.*\)\(\s*\/\*\)\(.*\)$/\1 #\2\3/; s/^\(\s*\*\)\(.*\)$/#\1\2/' extra/hip_gpu_driver/nvd.h >> $BASE/amd_gpu.py # comments
-  sed 's/^\(.*\)\(\s*\/\*\)\(.*\)$/\1 #\2\3/; s/^\(\s*\*\)\(.*\)$/#\1\2/' extra/hip_gpu_driver/sdma_v6_0_0_pkt_open.h >> $BASE/amd_gpu.py # comments
-  sed 's/^\(.*\)\(\s*\/\*\)\(.*\)$/\1 #\2\3/; s/^\(\s*\*\)\(.*\)$/#\1\2/' extra/hip_gpu_driver/gc_11_0_0_offset.h >> $BASE/amd_gpu.py # comments
-  sed 's/^\(.*\)\(\s*\/\*\)\(.*\)$/\1 #\2\3/; s/^\(\s*\*\)\(.*\)$/#\1\2/' extra/hip_gpu_driver/gc_10_3_0_offset.h >> $BASE/amd_gpu.py # comments
-  sed -i 's/^\/\//#/' $BASE/amd_gpu.py # // -> #
-  sed -i 's/#\s*define\s*\([^ \t]*\)(\([^)]*\))\s*\(.*\)/def \1(\2): return \3/' $BASE/amd_gpu.py # #define name(x) (smth) -> def name(x): return (smth)
-  sed -i '/#\s*define\s\+\([^ \t]\+\)\s\+\([^ ]\+\)/s//\1 = \2/' $BASE/amd_gpu.py # #define name val -> name = val
-
-  # sed -e '/^reg/s/^\(reg[^ ]*\) [^ ]* \([^ ]*\) .*/\1 = \2/' \
-  #   -e '/^ix/s/^\(ix[^ ]*\) [^ ]* \([^ ]*\) .*/\1 = \2/' \
-  #   -e '/^[ \t]/d' \
-  #   extra/hip_gpu_driver/gc_11_0_0.reg >> $BASE/amd_gpu.py
-
  fixup $BASE/amd_gpu.py
  sed -i "s\import ctypes\import ctypes, os\g" $BASE/amd_gpu.py
  python3 -c "import tinygrad.runtime.autogen.amd_gpu"
@@ -252,28 +190,23 @@ generate_hsa() {
 }

 generate_io_uring() {
-  clang2py \
+  clang2py -k cdefstum \
    /usr/include/liburing.h \
    /usr/include/linux/io_uring.h \
    -o $BASE/io_uring.py

-  # clang2py can't parse defines
  sed -r '/^#define __NR_io_uring/ s/^#define __(NR_io_uring[^ ]+) (.*)$/\1 = \2/; t; d' /usr/include/asm-generic/unistd.h >> $BASE/io_uring.py # io_uring syscalls numbers
-  sed -r '/^#define\s+([^ \t]+)\s+([^ \t]+)/ s/^#define\s+([^ \t]+)\s*([^/]*).*$/\1 = \2/; s/1U/1/g; s/0ULL/0/g; t; d' /usr/include/linux/io_uring.h >> $BASE/io_uring.py # #define name (val) -> name = val
-
  fixup $BASE/io_uring.py
 }

 generate_libc() {
-  clang2py \
+  clang2py -k cdefstum \
    $(dpkg -L libc6-dev | grep sys/mman.h) \
    $(dpkg -L libc6-dev | grep sys/syscall.h) \
    /usr/include/elf.h \
    /usr/include/unistd.h \
    -o $BASE/libc.py

-  process_cdefines "/usr/include/elf.h" "$BASE/libc.py"
-
  sed -i "s\import ctypes\import ctypes, ctypes.util, os\g" $BASE/libc.py
  sed -i "s\FIXME_STUB\libc\g" $BASE/libc.py
  sed -i "s\FunctionFactoryStub()\ctypes.CDLL(ctypes.util.find_library('c'))\g" $BASE/libc.py
@@ -281,6 +214,35 @@ generate_libc() {
  fixup $BASE/libc.py
 }

+generate_kgsl() {
+  clang2py extra/qcom_gpu_driver/msm_kgsl.h -o $BASE/kgsl.py -k cdefstum
+  fixup $BASE/kgsl.py
+  sed -i "s\import ctypes\import ctypes, os\g" $BASE/kgsl.py
+  sed -nE 's/#define ([A-Za-z0-9_]+)_SHIFT\s*[^\S\r\n]*[0-9]*$/def \1(val): return (val << \1_SHIFT) \& \1_MASK/p' extra/qcom_gpu_driver/msm_kgsl.h >> $BASE/kgsl.py
+  python3 -c "import tinygrad.runtime.autogen.kgsl"
+}
+
+generate_adreno() {
+  clang2py extra/qcom_gpu_driver/a6xx.xml.h -o $BASE/adreno.py -k cestum
+  sed -nE 's/#define ([A-Za-z0-9_]+)__SHIFT\s*[^\S\r\n]*[0-9]*$/def \1(val): return (val << \1__SHIFT) \& \1__MASK/p' extra/qcom_gpu_driver/a6xx.xml.h >> $BASE/adreno.py
+  fixup $BASE/adreno.py
+  sed -i "s\import ctypes\import ctypes, os\g" $BASE/adreno.py
+  python3 -c "import tinygrad.runtime.autogen.adreno"
+}
+
+generate_qcom() {
+  clang2py -k cdefstum \
+    extra/dsp/include/ion.h \
+    extra/dsp/include/msm_ion.h \
+    extra/dsp/include/adsprpc_shared.h \
+    extra/dsp/include/remote_default.h \
+    extra/dsp/include/apps_std.h \
+    -o $BASE/qcom_dsp.py
+
+  fixup $BASE/qcom_dsp.py
+  python3 -c "import tinygrad.runtime.autogen.qcom_dsp"
+}
+
 if [ "$1" == "opencl" ]; then generate_opencl
 elif [ "$1" == "hip" ]; then generate_hip
 elif [ "$1" == "comgr" ]; then generate_comgr
@@ -290,8 +252,11 @@ elif [ "$1" == "hsa" ]; then generate_hsa
 elif [ "$1" == "kfd" ]; then generate_kfd
 elif [ "$1" == "nv" ]; then generate_nv
 elif [ "$1" == "amd" ]; then generate_amd
+elif [ "$1" == "qcom" ]; then generate_qcom
 elif [ "$1" == "io_uring" ]; then generate_io_uring
 elif [ "$1" == "libc" ]; then generate_libc
+elif [ "$1" == "kgsl" ]; then generate_kgsl
+elif [ "$1" == "adreno" ]; then generate_adreno
 elif [ "$1" == "all" ]; then generate_opencl; generate_hip; generate_comgr; generate_cuda; generate_nvrtc; generate_hsa; generate_kfd; generate_nv; generate_amd; generate_io_uring; generate_libc
 else echo "usage: $0 <type>"
 fi
@@ -37,9 +37,9 @@ print("******** second, the Device ***********")
 DEVICE = "CLANG"   # NOTE: you can change this!

 import struct
-from tinygrad.dtype import dtypes
+from tinygrad.dtype import PtrDType, dtypes
 from tinygrad.device import Buffer, Device
-from tinygrad.ops import LazyOp, BufferOps, MemBuffer, BinaryOps, MetaOps
+from tinygrad.ops import BinaryOps, MetaOps, UOp, UOps
 from tinygrad.shape.shapetracker import ShapeTracker

 # allocate some buffers + load in values
@@ -49,16 +49,18 @@ b = Buffer(DEVICE, 1, dtypes.int32).allocate().copyin(memoryview(bytearray(struc
 # NOTE: a._buf is the same as the return from MallocAllocator.alloc

 # describe the computation
-ld_1 = LazyOp(BufferOps.LOAD, (), MemBuffer(1, dtypes.int32, ShapeTracker.from_shape((1,))))
-ld_2 = LazyOp(BufferOps.LOAD, (), MemBuffer(2, dtypes.int32, ShapeTracker.from_shape((1,))))
-alu = LazyOp(BinaryOps.ADD, (ld_1, ld_2))
-st_0 = LazyOp(BufferOps.STORE, (alu,), MemBuffer(0, dtypes.int32, ShapeTracker.from_shape((1,))))
-k = LazyOp(MetaOps.KERNEL, (st_0,))
+buf_1 = UOp(UOps.DEFINE_GLOBAL, PtrDType(dtypes.int32), (), 1)
+buf_2 = UOp(UOps.DEFINE_GLOBAL, PtrDType(dtypes.int32), (), 2)
+ld_1 = UOp(UOps.LOAD, dtypes.int32, (buf_1, ShapeTracker.from_shape((1,)).to_uop()))
+ld_2 = UOp(UOps.LOAD, dtypes.int32, (buf_2, ShapeTracker.from_shape((1,)).to_uop()))
+alu = ld_1 + ld_2
+output_buf = UOp(UOps.DEFINE_GLOBAL, PtrDType(dtypes.int32), (), 0)
+st_0 = UOp(UOps.STORE, dtypes.void, (output_buf, ShapeTracker.from_shape((1,)).to_uop(), alu))
+s = UOp(UOps.SINK, dtypes.void, (st_0,))

 # convert the computation to a "linearized" format (print the format)
 from tinygrad.engine.realize import get_kernel, CompiledRunner
-kernel = get_kernel(Device[DEVICE].renderer, k).linearize()
-kernel.uops.print()
+kernel = get_kernel(Device[DEVICE].renderer, s).linearize()

 # compile a program (and print the source)
 fxn = CompiledRunner(kernel.to_program())
@@ -74,7 +76,7 @@ assert out.as_buffer().cast('I')[0] == 5

 print("******** third, the LazyBuffer ***********")

-from tinygrad.lazy import LazyBuffer
+from tinygrad.engine.lazy import LazyBuffer
 from tinygrad.engine.realize import run_schedule
 from tinygrad.engine.schedule import create_schedule

@@ -87,7 +89,7 @@ del a.srcs
 del b.srcs

 # describe the computation
-out = a.e(BinaryOps.ADD, b)
+out = a.alu(BinaryOps.ADD, b)

 # schedule the computation as a list of kernels
 sched = create_schedule([out])
@@ -11,7 +11,7 @@ There is a good [bunch of tutorials](https://mesozoic-egg.github.io/tinygrad-not

 Everything in [Tensor](../tensor/index.md) is syntactic sugar around [function.py](function.md), where the forwards and backwards passes are implemented for the different functions. There's about 25 of them, implemented using about 20 basic ops. Those basic ops go on to construct a graph of:

-::: tinygrad.lazy.LazyBuffer
+::: tinygrad.engine.lazy.LazyBuffer
    options:
        show_source: false

@@ -6,7 +6,7 @@ The main aspect of HCQ-compatible runtimes is how they interact with devices. In

 ### Command Queues

-To interact with devices, there are 2 types of queues: `HWComputeQueue` and `HWCopyQueue`. Commands which are defined in a base `HWCommandQueue` class should be supported by both queues. These methods are timestamp and synchronization methods like [signal](#tinygrad.device.HWCommandQueue.signal) and [wait](#tinygrad.device.HWCommandQueue.wait).
+To interact with devices, there are 2 types of queues: `HWComputeQueue` and `HWCopyQueue`. Commands which are defined in a base `HWCommandQueue` class should be supported by both queues. These methods are timestamp and synchronization methods like [signal](#tinygrad.runtime.support.hcq.HWCommandQueue.signal) and [wait](#tinygrad.runtime.support.hcq.HWCommandQueue.wait).

 For example, the following Python code enqueues a wait, execute, and signal command on the HCQ-compatible device:
 ```python
@@ -18,7 +18,7 @@ HWComputeQueue().wait(signal_to_wait, value_to_wait) \

 Each runtime should implement the required functions that are defined in the `HWCommandQueue`, `HWComputeQueue`, and `HWCopyQueue` classes.

-::: tinygrad.device.HWCommandQueue
+::: tinygrad.runtime.support.hcq.HWCommandQueue
    options:
        members: [
            "signal",
@@ -31,7 +31,7 @@ Each runtime should implement the required functions that are defined in the `HW
        ]
        show_source: false

-::: tinygrad.device.HWComputeQueue
+::: tinygrad.runtime.support.hcq.HWComputeQueue
    options:
        members: [
            "memory_barrier",
@@ -40,7 +40,7 @@ Each runtime should implement the required functions that are defined in the `HW
        ]
        show_source: false

-::: tinygrad.device.HWCopyQueue
+::: tinygrad.runtime.support.hcq.HWCopyQueue
    options:
        members: [
            "copy",
@@ -52,7 +52,7 @@ Each runtime should implement the required functions that are defined in the `HW

 To implement custom commands in the queue, use the @hcq_command decorator for your command implementations.

-::: tinygrad.device.hcq_command
+::: tinygrad.runtime.support.hcq.hcq_command
    options:
        members: [
            "copy",
@@ -64,7 +64,7 @@ To implement custom commands in the queue, use the @hcq_command decorator for yo

 The `HCQCompiled` class defines the API for HCQ-compatible devices. This class serves as an abstract base class that device-specific implementations should inherit from and implement.

-::: tinygrad.device.HCQCompiled
+::: tinygrad.runtime.support.hcq.HCQCompiled
    options:
        show_source: false

@@ -72,7 +72,7 @@ The `HCQCompiled` class defines the API for HCQ-compatible devices. This class s

 Signals are device-dependent structures used for synchronization and timing in HCQ-compatible devices. They should be designed to record both a `value` and a `timestamp` within the same signal. HCQ-compatible backend implementations should use `HCQSignal` as a base class.

-::: tinygrad.device.HCQSignal
+::: tinygrad.runtime.support.hcq.HCQSignal
    options:
        members: [value, timestamp, wait]
        show_source: false
@@ -99,7 +99,7 @@ Each HCQ-compatible device must allocate two signals for global synchronization

 The `HCQAllocator` base class simplifies allocator logic by leveraging [command queues](#command-queues) abstractions. This class efficiently handles copy and transfer operations, leaving only the alloc and free functions to be implemented by individual backends.

-::: tinygrad.device.HCQAllocator
+::: tinygrad.runtime.support.hcq.HCQAllocator
    options:
        members: [
            "_alloc",
@@ -111,7 +111,7 @@ The `HCQAllocator` base class simplifies allocator logic by leveraging [command

 Backends must adhere to the `HCQBuffer` protocol when returning allocation results.

-::: tinygrad.device.HCQBuffer
+::: tinygrad.runtime.support.hcq.HCQBuffer
    options:
        members: true
        show_source: false
@@ -120,7 +120,7 @@ Backends must adhere to the `HCQBuffer` protocol when returning allocation resul

 `HCQProgram` is a base class for defining programs compatible with HCQ-enabled devices. It provides a flexible framework for handling different argument layouts (see `HCQArgsState`).

-::: tinygrad.device.HCQProgram
+::: tinygrad.runtime.support.hcq.HCQProgram
    options:
        members: true
        show_source: false
@@ -129,11 +129,13 @@ Backends must adhere to the `HCQBuffer` protocol when returning allocation resul

 `HCQArgsState` is a base class for managing the argument state for HCQ programs. Backend implementations should create a subclass of `HCQArgsState` to manage arguments for the given program.

-::: tinygrad.device.HCQArgsState
+::: tinygrad.runtime.support.hcq.HCQArgsState
    options:
        members: true
        show_source: false

+**Lifetime**: The `HCQArgsState` is passed to `HWComputeQueue.exec` and is guaranteed not to be freed until `HWComputeQueue.submit` for the same queue is called.
+
 ### Synchronization

 HCQ-compatible devices use a global timeline signal for synchronizing all operations. This mechanism ensures proper ordering and completion of tasks across the device. By convention, `self.timeline_value` points to the next value to signal. So, to wait for all previous operations on the device to complete, wait for `self.timeline_value - 1` value. The following Python code demonstrates the typical usage of signals to synchronize execution to other operations on the device:
@@ -0,0 +1,11 @@
+::: tinygrad.ops.UOp
+    options:
+        members: false
+        members_order: source
+        show_labels: false
+
+::: tinygrad.ops.UOps
+    options:
+        members: true
+        members_order: source
+        show_labels: false
@@ -10,7 +10,9 @@ cd tinygrad
 python3 -m pip install -e .
 ```

-After you have installed tinygrad, try the [MNIST tutorial](mnist.md)
+After you have installed tinygrad, try the [MNIST tutorial](mnist.md).
+
+If you are new to tensor libraries, learn how to use them by solving puzzles from [tinygrad-tensor-puzzles](https://github.com/obadakhalili/tinygrad-tensor-puzzles).

 We also have [developer docs](developer/developer.md), and Di Zhu has created a [bunch of tutorials](https://mesozoic-egg.github.io/tinygrad-notes/) to help understand how tinygrad works.

@@ -98,6 +98,14 @@ timeit.repeat(step, repeat=5, number=1)

 So around 75 ms on T4 colab.

+If you want to see a breakdown of the time by kernel:
+
+```python
+from tinygrad import GlobalCounters, Context
+GlobalCounters.reset()
+with Context(DEBUG=2): step()
+```
+
 ### Why so slow?

 Unlike PyTorch, tinygrad isn't designed to be fast like that. While 75 ms for one step is plenty fast for debugging, it's not great for training. Here, we introduce the first quintessentially tinygrad concept, the `TinyJit`.
@@ -12,6 +12,7 @@
 ::: tinygrad.nn.LayerNorm2d
 ::: tinygrad.nn.RMSNorm
 ::: tinygrad.nn.Embedding
+::: tinygrad.nn.LSTMCell

 ## Optimizers

@@ -5,7 +5,8 @@ tinygrad supports various runtimes, enabling your code to scale across a wide ra
 | Runtime | Description | Requirements |
 |---------|-------------|--------------|
 | [NV](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_nv.py) | Provides acceleration for NVIDIA GPUs | Ampere/Ada series GPUs |
-| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | RDNA3 series GPUs |
+| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | RDNA2/RDNA3 series GPUs |
+| [QCOM](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_qcom.py) | Provides acceleration for QCOM GPUs | 6xx series GPUs |
 | [METAL](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_metal.py) | Utilizes Metal for acceleration on Apple devices | M1+ Macs; Metal 3.0+ for `bfloat` support |
 | [CUDA](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cuda.py) | Utilizes CUDA for acceleration on NVIDIA GPUs | NVIDIA GPU with CUDA support |
 | [GPU (OpenCL)](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_gpu.py) | Accelerates computations using OpenCL on GPUs | OpenCL 2.0 compatible device |
@@ -9,6 +9,7 @@
 ::: tinygrad.Tensor.full_like
 ::: tinygrad.Tensor.zeros_like
 ::: tinygrad.Tensor.ones_like
+::: tinygrad.Tensor.from_blob

 ## Creation (random)

@@ -1,6 +1,7 @@
 ## Reduce

 ::: tinygrad.Tensor.sum
+::: tinygrad.Tensor.prod
 ::: tinygrad.Tensor.max
 ::: tinygrad.Tensor.min
 ::: tinygrad.Tensor.any
@@ -8,9 +9,11 @@
 ::: tinygrad.Tensor.mean
 ::: tinygrad.Tensor.var
 ::: tinygrad.Tensor.std
+::: tinygrad.Tensor.std_mean
 ::: tinygrad.Tensor.softmax
 ::: tinygrad.Tensor.log_softmax
 ::: tinygrad.Tensor.logsumexp
+::: tinygrad.Tensor.logcumsumexp
 ::: tinygrad.Tensor.argmax
 ::: tinygrad.Tensor.argmin

@@ -40,3 +43,4 @@
 ::: tinygrad.Tensor.binary_crossentropy
 ::: tinygrad.Tensor.binary_crossentropy_logits
 ::: tinygrad.Tensor.sparse_categorical_crossentropy
+::: tinygrad.Tensor.cross_entropy
@@ -12,7 +12,7 @@ We don't have a stupid cloud service, you don't have to create a tiny account to

 ## Plugging it in

-tinybox has two 1600W PSUs, which together exceed the capacity of most 120V household circuits. Fortunately, it comes with two plugs. You'll want to plug each plug into a different circuit. You can verify that they are different circuits by flipping the breaker and seeing what turns off. If you have at least a 120V 30A or 220V 15A circuit, you are welcome to use only that one.
+tinybox has two 1600W PSUs, which together exceed the capacity of most 120V household circuits. Fortunately, it comes with two plugs. You'll want to plug each plug into a different circuit. You can verify that they are different circuits by flipping the breaker and seeing what turns off. If you have at least a 120V 30A or 220V 20A circuit, you are welcome to use only that one.

 You'll also want to connect the Ethernet port without a rubber stopper to your home network.

@@ -32,17 +32,16 @@ tinybox ships with a relatively basic install of Ubuntu 22.04. To do initial set

 `ipmitool -H <BMC IP> -U admin -P <BMC PW> -I lanplus sol activate`

-The default username is `tiny` and the default password is `tiny`. Once you are logged in, you can add an SSH key to authorized keys to connect over SSH (on the normal IP). Exit `ipmitool` with `~~.`
+The default username is `tiny` and the default password is `tiny`. Once you are logged in, you can add an SSH key to authorized keys to connect over SSH (on the normal IP). Exit `ipmitool` with `~.` after a newline.

 The BMC also has a web interface you can use if you find that easier.

 ## Changing the BMC password

-If you try to change the BMC password over IPMI or over the web interface, you will notice that it does not persist across reboots, and the password will revert to the one displayed on the screen.
+It is recommended that you change the BMC password after setting up the box, as the password on the screen is only the initial password.

-If you want to change the password imperatively, remove the `/root/.bmc_password` file and then set the password, the BMC password will also no longer be displayed on the screen. Additionally, you may modify the password stored in the `/root/.bmc_password` file to one that you choose if you still want it displayed on the screen.
-
-Reboot after making these changes.
+If you do decide to change the BMC password and no longer want the initial password to be displayed, remove the `/root/.bmc_password` file.
+Reboot after making these changes or restart the `displayservice.service` service.

 ## What do I use it for?

@@ -0,0 +1,168 @@
+import time
+start_tm = time.perf_counter()
+import math
+from typing import Tuple, cast
+import numpy as np
+from tinygrad import Tensor, nn, GlobalCounters, TinyJit, dtypes
+from tinygrad.helpers import partition, trange, getenv, Context
+from extra.lr_scheduler import OneCycleLR
+
+dtypes.default_float = dtypes.half
+
+# from https://github.com/tysam-code/hlb-CIFAR10/blob/main/main.py
+batchsize = getenv("BS", 1024)
+bias_scaler = 64
+hyp = {
+  'opt': {
+    'bias_lr':        1.525 * bias_scaler/512, # TODO: Is there maybe a better way to express the bias and batchnorm scaling? :'))))
+    'non_bias_lr':    1.525 / 512,
+    'bias_decay':     6.687e-4 * batchsize/bias_scaler,
+    'non_bias_decay': 6.687e-4 * batchsize,
+    'scaling_factor': 1./9,
+    'percent_start': .23,
+    'loss_scale_scaler': 1./32, # * Regularizer inside the loss summing (range: ~1/512 - 16+). FP8 should help with this somewhat too, whenever it comes out. :)
+  },
+  'net': {
+    'whitening': {
+      'kernel_size': 2,
+      'num_examples': 50000,
+    },
+    'batch_norm_momentum': .4, # * Don't forget momentum is 1 - momentum here (due to a quirk in the original paper... >:( )
+    'cutmix_size': 3,
+    'cutmix_epochs': 6,
+    'pad_amount': 2,
+    'base_depth': 64 ## This should be a factor of 8 in some way to stay tensor core friendly
+  },
+  'misc': {
+    'ema': {
+      'epochs': 10, # Slight bug in that this counts only full epochs and then additionally runs the EMA for any fractional epochs at the end too
+      'decay_base': .95,
+      'decay_pow': 3.,
+      'every_n_steps': 5,
+    },
+    'train_epochs': 12,
+    #'train_epochs': 12.1,
+    'device': 'cuda',
+    'data_location': 'data.pt',
+  }
+}
+
+scaler = 2. ## You can play with this on your own if you want, for the first beta I wanted to keep things simple (for now) and leave it out of the hyperparams dict
+depths = {
+  'init':   round(scaler**-1*hyp['net']['base_depth']), # 32  w/ scaler at base value
+  'block1': round(scaler** 0*hyp['net']['base_depth']), # 64  w/ scaler at base value
+  'block2': round(scaler** 2*hyp['net']['base_depth']), # 256 w/ scaler at base value
+  'block3': round(scaler** 3*hyp['net']['base_depth']), # 512 w/ scaler at base value
+  'num_classes': 10
+}
+whiten_conv_depth = 3*hyp['net']['whitening']['kernel_size']**2
+
+class ConvGroup:
+  def __init__(self, channels_in, channels_out):
+    self.conv1 = nn.Conv2d(channels_in, channels_out, kernel_size=3, padding=1, bias=False)
+    self.conv2 = nn.Conv2d(channels_out, channels_out, kernel_size=3, padding=1, bias=False)
+    self.norm1 = nn.BatchNorm(channels_out, track_running_stats=False, eps=1e-12, momentum=hyp['net']['batch_norm_momentum'])
+    self.norm2 = nn.BatchNorm(channels_out, track_running_stats=False, eps=1e-12, momentum=hyp['net']['batch_norm_momentum'])
+    cast(Tensor, self.norm1.weight).requires_grad = False
+    cast(Tensor, self.norm2.weight).requires_grad = False
+  def __call__(self, x:Tensor) -> Tensor:
+    x =    self.norm1(self.conv1(x).max_pool2d().float()).cast(dtypes.default_float).quick_gelu()
+    return self.norm2(self.conv2(x).float()).cast(dtypes.default_float).quick_gelu()
+
+class SpeedyConvNet:
+  def __init__(self):
+    self.whiten = nn.Conv2d(3, 2*whiten_conv_depth, kernel_size=hyp['net']['whitening']['kernel_size'], padding=0, bias=False)
+    self.conv_group_1 = ConvGroup(2*whiten_conv_depth, depths['block1'])
+    self.conv_group_2 = ConvGroup(depths['block1'], depths['block2'])
+    self.conv_group_3 = ConvGroup(depths['block2'], depths['block3'])
+    self.linear = nn.Linear(depths['block3'], depths['num_classes'], bias=False)
+  def __call__(self, x:Tensor) -> Tensor:
+    x = self.whiten(x).quick_gelu()
+    x = x.sequential([self.conv_group_1, self.conv_group_2, self.conv_group_3])
+    return self.linear(x.max(axis=(2,3))) * hyp['opt']['scaling_factor']
+
+if __name__ == "__main__":
+  # *** dataset ***
+  X_train, Y_train, X_test, Y_test = nn.datasets.cifar()
+  # TODO: without this line indexing doesn't fuse!
+  X_train, Y_train, X_test, Y_test = [x.contiguous() for x in [X_train, Y_train, X_test, Y_test]]
+  cifar10_std, cifar10_mean = X_train.float().std_mean(axis=(0, 2, 3))
+  def preprocess(X:Tensor, Y:Tensor) -> Tuple[Tensor, Tensor]:
+    return ((X - cifar10_mean.view(1, -1, 1, 1)) / cifar10_std.view(1, -1, 1, 1)).cast(dtypes.default_float), Y.one_hot(depths['num_classes'])
+
+  # *** model ***
+  model = SpeedyConvNet()
+  state_dict = nn.state.get_state_dict(model)
+
+  #for k,v in nn.state.torch_load("/tmp/cifar_net.pt").items(): print(k)
+
+  params_bias, params_non_bias = partition(state_dict.items(), lambda x: 'bias' in x[0])
+  opt_bias     = nn.optim.SGD([x[1] for x in params_bias],     lr=0.01, momentum=.85, nesterov=True, weight_decay=hyp['opt']['bias_decay'])
+  opt_non_bias = nn.optim.SGD([x[1] for x in params_non_bias], lr=0.01, momentum=.85, nesterov=True, weight_decay=hyp['opt']['non_bias_decay'])
+  opt = nn.optim.OptimizerGroup(opt_bias, opt_non_bias)
+
+  num_steps_per_epoch      = X_train.size(0) // batchsize
+  total_train_steps        = math.ceil(num_steps_per_epoch * hyp['misc']['train_epochs'])
+  loss_batchsize_scaler = 512/batchsize
+
+  pct_start = hyp['opt']['percent_start']
+  initial_div_factor = 1e16 # basically to make the initial lr ~0 or so :D
+  final_lr_ratio = .07 # Actually pretty important, apparently!
+  lr_sched_bias     = OneCycleLR(opt_bias,     max_lr=hyp['opt']['bias_lr'],     pct_start=pct_start, div_factor=initial_div_factor, final_div_factor=1./(initial_div_factor*final_lr_ratio), total_steps=total_train_steps)
+  lr_sched_non_bias = OneCycleLR(opt_non_bias, max_lr=hyp['opt']['non_bias_lr'], pct_start=pct_start, div_factor=initial_div_factor, final_div_factor=1./(initial_div_factor*final_lr_ratio), total_steps=total_train_steps)
+
+  def loss_fn(out, Y):
+    return out.cross_entropy(Y, reduction='none', label_smoothing=0.2).mul(hyp['opt']['loss_scale_scaler']*loss_batchsize_scaler).sum().div(hyp['opt']['loss_scale_scaler'])
+
+  @TinyJit
+  @Tensor.train()
+  def train_step(idxs:Tensor) -> Tensor:
+    with Context(SPLIT_REDUCEOP=0, FUSE_ARANGE=1):
+      X = X_train[idxs]
+      Y = Y_train[idxs].realize(X)
+    X, Y = preprocess(X, Y)
+    out = model(X)
+    loss = loss_fn(out, Y)
+    opt.zero_grad()
+    loss.backward()
+    opt.step()
+    lr_sched_bias.step()
+    lr_sched_non_bias.step()
+    return loss / (batchsize*loss_batchsize_scaler)
+
+  eval_batchsize = 2500
+  @TinyJit
+  @Tensor.test()
+  def val_step() -> Tuple[Tensor, Tensor]:
+    # TODO with Tensor.no_grad()
+    Tensor.no_grad = True
+    loss, acc = [], []
+    for i in range(0, X_test.size(0), eval_batchsize):
+      X, Y = preprocess(X_test[i:i+eval_batchsize], Y_test[i:i+eval_batchsize])
+      out = model(X)
+      loss.append(loss_fn(out, Y))
+      acc.append((out.argmax(-1).one_hot(depths['num_classes']) * Y).sum() / eval_batchsize)
+    ret = Tensor.stack(*loss).mean() / (batchsize*loss_batchsize_scaler), Tensor.stack(*acc).mean()
+    Tensor.no_grad = False
+    return ret
+
+  np.random.seed(1337)
+  for epoch in range(math.ceil(hyp['misc']['train_epochs'])):
+    # TODO: move to tinygrad
+    gst = time.perf_counter()
+    idxs = np.arange(X_train.shape[0])
+    np.random.shuffle(idxs)
+    tidxs = Tensor(idxs, dtype='int')[:num_steps_per_epoch*batchsize].reshape(num_steps_per_epoch, batchsize)  # NOTE: long doesn't fold
+    train_loss:float = 0
+    for epoch_step in (t:=trange(num_steps_per_epoch)):
+      st = time.perf_counter()
+      GlobalCounters.reset()
+      loss = train_step(tidxs[epoch_step].contiguous()).float().item()
+      t.set_description(f"*** loss: {loss:5.3f}   lr: {opt_non_bias.lr.item():.6f}"
+                         f"   tm: {(et:=(time.perf_counter()-st))*1000:6.2f} ms {GlobalCounters.global_ops/(1e9*et):7.0f} GFLOPS")
+      train_loss += loss
+    gmt = time.perf_counter()
+    GlobalCounters.reset()
+    val_loss, acc = [x.float().item() for x in val_step()]
+    get = time.perf_counter()
+    print(f"\033[F*** epoch {epoch:3d}       tm: {(gmt-gst):5.2f} s    val_tm: {(get-gmt):5.2f} s   train_loss: {train_loss/num_steps_per_epoch:5.3f}   val_loss: {val_loss:5.3f}   eval acc: {acc*100:5.2f}%  @ {get-start_tm:6.2f} s  ")
@@ -18,22 +18,23 @@ class Model:
  def __call__(self, x:Tensor) -> Tensor: return x.sequential(self.layers)

 if __name__ == "__main__":
-  X_train, Y_train, X_test, Y_test = mnist()
+  X_train, Y_train, X_test, Y_test = mnist(fashion=getenv("FASHION"))

  model = Model()
  opt = nn.optim.Adam(nn.state.get_parameters(model))

  @TinyJit
+  @Tensor.train()
  def train_step() -> Tensor:
-    with Tensor.train():
-      opt.zero_grad()
-      samples = Tensor.randint(getenv("BS", 512), high=X_train.shape[0])
-      # TODO: this "gather" of samples is very slow. will be under 5s when this is fixed
-      loss = model(X_train[samples]).sparse_categorical_crossentropy(Y_train[samples]).backward()
-      opt.step()
-      return loss
+    opt.zero_grad()
+    samples = Tensor.randint(getenv("BS", 512), high=X_train.shape[0])
+    # TODO: this "gather" of samples is very slow. will be under 5s when this is fixed
+    loss = model(X_train[samples]).sparse_categorical_crossentropy(Y_train[samples]).backward()
+    opt.step()
+    return loss

  @TinyJit
+  @Tensor.test()
  def get_test_acc() -> Tensor: return (model(X_test).argmax(axis=1) == Y_test).mean()*100

  test_acc = float('nan')
@@ -46,4 +47,4 @@ if __name__ == "__main__":
  # verify eval acc
  if target := getenv("TARGET_EVAL_ACC_PCT", 0.0):
    if test_acc >= target and test_acc != 100.0: print(colored(f"{test_acc=} >= {target}", "green"))
-    else: raise ValueError(colored(f"{test_acc=} < {target}", "red"))
+    else: raise ValueError(colored(f"{test_acc=} < {target}", "red"))
@@ -4,7 +4,7 @@ from tinygrad import Tensor, TinyJit, nn, GlobalCounters, Device
 from tinygrad.helpers import getenv, colored, trange
 from tinygrad.nn.datasets import mnist

-GPUS = [f'{Device.DEFAULT}:{i}' for i in range(getenv("GPUS", 2))]
+GPUS = tuple(f'{Device.DEFAULT}:{i}' for i in range(getenv("GPUS", 2)))

 class Model:
  def __init__(self):
@@ -33,7 +33,7 @@ if __name__ == "__main__":
  def train_step() -> Tensor:
    with Tensor.train():
      opt.zero_grad()
-      samples = Tensor.randint(512, high=X_train.shape[0])
+      samples = Tensor.randint(getenv("BS", 512), high=X_train.shape[0])
      Xt, Yt = X_train[samples].shard_(GPUS, axis=0), Y_train[samples].shard_(GPUS, axis=0)  # we shard the data on axis 0
      # TODO: this "gather" of samples is very slow. will be under 5s when this is fixed
      loss = model(Xt).sparse_categorical_crossentropy(Yt).backward()
@@ -44,7 +44,7 @@ if __name__ == "__main__":
  def get_test_acc() -> Tensor: return (model(X_test).argmax(axis=1) == Y_test).mean()*100

  test_acc = float('nan')
-  for i in (t:=trange(70)):
+  for i in (t:=trange(getenv("STEPS", 70))):
    GlobalCounters.reset()   # NOTE: this makes it nice for DEBUG=2 timing
    loss = train_step()
    if i%10 == 9: test_acc = get_test_acc().item()
@@ -53,4 +53,4 @@ if __name__ == "__main__":
  # verify eval acc
  if target := getenv("TARGET_EVAL_ACC_PCT", 0.0):
    if test_acc >= target: print(colored(f"{test_acc=} >= {target}", "green"))
-    else: raise ValueError(colored(f"{test_acc=} < {target}", "red"))
+    else: raise ValueError(colored(f"{test_acc=} < {target}", "red"))
@@ -0,0 +1,496 @@
+# pip3 install sentencepiece
+
+# This file incorporates code from the following:
+# Github Name                    | License | Link
+# black-forest-labs/flux         | Apache  | https://github.com/black-forest-labs/flux/tree/main/model_licenses
+
+from tinygrad import Tensor, nn, dtypes, TinyJit
+from tinygrad.nn.state import safe_load, load_state_dict
+from tinygrad.helpers import fetch, tqdm, colored
+from sdxl import FirstStage
+from extra.models.clip import FrozenClosedClipEmbedder
+from extra.models.t5 import T5Embedder
+import numpy as np
+
+import math, time, argparse, tempfile
+from typing import List, Dict, Optional, Union, Tuple, Callable
+from dataclasses import dataclass
+from pathlib import Path
+from PIL import Image
+
+urls:dict = {
+  "flux-schnell": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/flux1-schnell.safetensors",
+  "flux-dev": "https://huggingface.co/camenduru/FLUX.1-dev/resolve/main/flux1-dev.sft",
+  "ae": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/ae.safetensors",
+  "T5_1_of_2": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/text_encoder_2/model-00001-of-00002.safetensors",
+  "T5_2_of_2": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/text_encoder_2/model-00002-of-00002.safetensors",
+  "T5_tokenizer": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/tokenizer_2/spiece.model",
+  "clip": "https://huggingface.co/black-forest-labs/FLUX.1-schnell/resolve/main/text_encoder/model.safetensors"
+}
+
+def tensor_identity(x:Tensor) -> Tensor: return x
+
+class AutoEncoder:
+  def __init__(self, scale_factor:float, shift_factor:float):
+    self.decoder = FirstStage.Decoder(128, 3, 3, 16, [1, 2, 4, 4], 2, 256)
+    self.scale_factor = scale_factor
+    self.shift_factor = shift_factor
+
+  def decode(self, z:Tensor) -> Tensor:
+    z = z / self.scale_factor + self.shift_factor
+    return self.decoder(z)
+
+# Conditioner
+class ClipEmbedder(FrozenClosedClipEmbedder):
+  def __call__(self, texts:Union[str, List[str], Tensor]) -> Tensor:
+    if isinstance(texts, str): texts = [texts]
+    assert isinstance(texts, (list,tuple)), f"expected list of strings, got {type(texts).__name__}"
+    tokens = Tensor.cat(*[Tensor(self.tokenizer.encode(text)) for text in texts], dim=0)
+    return self.transformer.text_model(tokens.reshape(len(texts),-1))[:, tokens.argmax(-1)]
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/math.py
+def attention(q:Tensor, k:Tensor, v:Tensor, pe:Tensor) -> Tensor:
+  q, k = apply_rope(q, k, pe)
+  x = Tensor.scaled_dot_product_attention(q, k, v)
+  return x.rearrange("B H L D -> B L (H D)")
+
+def rope(pos:Tensor, dim:int, theta:int) -> Tensor:
+  assert dim % 2 == 0
+  scale = Tensor.arange(0, dim, 2, dtype=dtypes.float32, device=pos.device) / dim # NOTE: this is torch.float64 in reference implementation
+  omega = 1.0 / (theta**scale)
+  out = Tensor.einsum("...n,d->...nd", pos, omega)
+  out = Tensor.stack(Tensor.cos(out), -Tensor.sin(out), Tensor.sin(out), Tensor.cos(out), dim=-1)
+  out = out.rearrange("b n d (i j) -> b n d i j", i=2, j=2)
+  return out.float()
+
+def apply_rope(xq:Tensor, xk:Tensor, freqs_cis:Tensor) -> Tuple[Tensor, Tensor]:
+  xq_ = xq.float().reshape(*xq.shape[:-1], -1, 1, 2)
+  xk_ = xk.float().reshape(*xk.shape[:-1], -1, 1, 2)
+  xq_out = freqs_cis[..., 0] * xq_[..., 0] + freqs_cis[..., 1] * xq_[..., 1]
+  xk_out = freqs_cis[..., 0] * xk_[..., 0] + freqs_cis[..., 1] * xk_[..., 1]
+  return xq_out.reshape(*xq.shape).cast(xq.dtype), xk_out.reshape(*xk.shape).cast(xk.dtype)
+
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/modules/layers.py
+class EmbedND:
+  def __init__(self, dim:int, theta:int, axes_dim:List[int]):
+    self.dim = dim
+    self.theta = theta
+    self.axes_dim = axes_dim
+
+  def __call__(self, ids:Tensor) -> Tensor:
+    n_axes = ids.shape[-1]
+    emb = Tensor.cat(*[rope(ids[..., i], self.axes_dim[i], self.theta) for i in range(n_axes)], dim=-3)
+    return emb.unsqueeze(1)
+
+class MLPEmbedder:
+  def __init__(self, in_dim:int, hidden_dim:int):
+    self.in_layer = nn.Linear(in_dim, hidden_dim, bias=True)
+    self.out_layer = nn.Linear(hidden_dim, hidden_dim, bias=True)
+
+  def __call__(self, x:Tensor) -> Tensor:
+    return self.out_layer(self.in_layer(x).silu())
+
+class QKNorm:
+  def __init__(self, dim:int):
+    self.query_norm = nn.RMSNorm(dim)
+    self.key_norm = nn.RMSNorm(dim)
+
+  def __call__(self, q:Tensor, k:Tensor) -> Tuple[Tensor, Tensor]:
+    return self.query_norm(q), self.key_norm(k)
+
+class SelfAttention:
+  def __init__(self, dim:int, num_heads:int = 8, qkv_bias:bool = False):
+    self.num_heads = num_heads
+    head_dim = dim // num_heads
+
+    self.qkv = nn.Linear(dim, dim * 3, bias=qkv_bias)
+    self.norm = QKNorm(head_dim)
+    self.proj = nn.Linear(dim, dim)
+
+  def __call__(self, x:Tensor, pe:Tensor) -> Tensor:
+    qkv = self.qkv(x)
+    q, k, v = qkv.rearrange("B L (K H D) -> K B H L D", K=3, H=self.num_heads)
+    q, k = self.norm(q, k)
+    x = attention(q, k, v, pe=pe)
+    return self.proj(x)
+
+@dataclass
+class ModulationOut:
+  shift:Tensor
+  scale:Tensor
+  gate:Tensor
+
+class Modulation:
+  def __init__(self, dim:int, double:bool):
+    self.is_double = double
+    self.multiplier = 6 if double else 3
+    self.lin = nn.Linear(dim, self.multiplier * dim, bias=True)
+
+  def __call__(self, vec:Tensor) -> Tuple[ModulationOut, Optional[ModulationOut]]:
+    out = self.lin(vec.silu())[:, None, :].chunk(self.multiplier, dim=-1)
+    return ModulationOut(*out[:3]), ModulationOut(*out[3:]) if self.is_double else None
+
+class DoubleStreamBlock:
+  def __init__(self, hidden_size:int, num_heads:int, mlp_ratio:float, qkv_bias:bool = False):
+    mlp_hidden_dim = int(hidden_size * mlp_ratio)
+    self.num_heads = num_heads
+    self.hidden_size = hidden_size
+    self.img_mod = Modulation(hidden_size, double=True)
+    self.img_norm1 = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+    self.img_attn = SelfAttention(dim=hidden_size, num_heads=num_heads, qkv_bias=qkv_bias)
+
+    self.img_norm2 = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+    self.img_mlp = [nn.Linear(hidden_size, mlp_hidden_dim, bias=True), Tensor.gelu, nn.Linear(mlp_hidden_dim, hidden_size, bias=True)]
+
+    self.txt_mod = Modulation(hidden_size, double=True)
+    self.txt_norm1 = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+    self.txt_attn = SelfAttention(dim=hidden_size, num_heads=num_heads, qkv_bias=qkv_bias)
+
+    self.txt_norm2 = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+    self.txt_mlp = [nn.Linear(hidden_size, mlp_hidden_dim, bias=True), Tensor.gelu, nn.Linear(mlp_hidden_dim, hidden_size, bias=True)]
+
+  def __call__(self, img:Tensor, txt:Tensor, vec:Tensor, pe:Tensor) -> tuple[Tensor, Tensor]:
+    img_mod1, img_mod2 = self.img_mod(vec)
+    txt_mod1, txt_mod2 = self.txt_mod(vec)
+    assert img_mod2 is not None and txt_mod2 is not None
+    # prepare image for attention
+    img_modulated = self.img_norm1(img)
+    img_modulated = (1 + img_mod1.scale) * img_modulated + img_mod1.shift
+    img_qkv = self.img_attn.qkv(img_modulated)
+    img_q, img_k, img_v = img_qkv.rearrange("B L (K H D) -> K B H L D", K=3, H=self.num_heads)
+    img_q, img_k = self.img_attn.norm(img_q, img_k)
+
+    # prepare txt for attention
+    txt_modulated = self.txt_norm1(txt)
+    txt_modulated = (1 + txt_mod1.scale) * txt_modulated + txt_mod1.shift
+    txt_qkv = self.txt_attn.qkv(txt_modulated)
+    txt_q, txt_k, txt_v = txt_qkv.rearrange("B L (K H D) -> K B H L D", K=3, H=self.num_heads)
+    txt_q, txt_k = self.txt_attn.norm(txt_q, txt_k)
+
+    # run actual attention
+    q = Tensor.cat(txt_q, img_q, dim=2)
+    k = Tensor.cat(txt_k, img_k, dim=2)
+    v = Tensor.cat(txt_v, img_v, dim=2)
+
+    attn = attention(q, k, v, pe=pe)
+    txt_attn, img_attn = attn[:, : txt.shape[1]], attn[:, txt.shape[1] :]
+
+    # calculate the img bloks
+    img = img + img_mod1.gate * self.img_attn.proj(img_attn)
+    img = img + img_mod2.gate * ((1 + img_mod2.scale) * self.img_norm2(img) + img_mod2.shift).sequential(self.img_mlp)
+
+    # calculate the txt bloks
+    txt = txt + txt_mod1.gate * self.txt_attn.proj(txt_attn)
+    txt = txt + txt_mod2.gate * ((1 + txt_mod2.scale) * self.txt_norm2(txt) + txt_mod2.shift).sequential(self.txt_mlp)
+    return img, txt
+
+
+class SingleStreamBlock:
+  """
+  A DiT block with parallel linear layers as described in
+  https://arxiv.org/abs/2302.05442 and adapted modulation interface.
+  """
+
+  def __init__(self,hidden_size:int, num_heads:int, mlp_ratio:float=4.0, qk_scale:Optional[float]=None):
+    self.hidden_dim = hidden_size
+    self.num_heads = num_heads
+    head_dim = hidden_size // num_heads
+    self.scale = qk_scale or head_dim**-0.5
+
+    self.mlp_hidden_dim = int(hidden_size * mlp_ratio)
+    # qkv and mlp_in
+    self.linear1 = nn.Linear(hidden_size, hidden_size * 3 + self.mlp_hidden_dim)
+    # proj and mlp_out
+    self.linear2 = nn.Linear(hidden_size + self.mlp_hidden_dim, hidden_size)
+
+    self.norm = QKNorm(head_dim)
+
+    self.hidden_size = hidden_size
+    self.pre_norm = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+
+    self.mlp_act = Tensor.gelu
+    self.modulation = Modulation(hidden_size, double=False)
+
+  def __call__(self, x:Tensor, vec:Tensor, pe:Tensor) -> Tensor:
+    mod, _ = self.modulation(vec)
+    x_mod = (1 + mod.scale) * self.pre_norm(x) + mod.shift
+    qkv, mlp = Tensor.split(self.linear1(x_mod), [3 * self.hidden_size, self.mlp_hidden_dim], dim=-1)
+    q, k, v = qkv.rearrange("B L (K H D) -> K B H L D", K=3, H=self.num_heads)
+    q, k = self.norm(q, k)
+
+    # compute attention
+    attn = attention(q, k, v, pe=pe)
+    # compute activation in mlp stream, cat again and run second linear layer
+    output = self.linear2(Tensor.cat(attn, self.mlp_act(mlp), dim=2))
+    return x + mod.gate * output
+
+
+class LastLayer:
+  def __init__(self, hidden_size:int, patch_size:int, out_channels:int):
+    self.norm_final = nn.LayerNorm(hidden_size, elementwise_affine=False, eps=1e-6)
+    self.linear = nn.Linear(hidden_size, patch_size * patch_size * out_channels, bias=True)
+    self.adaLN_modulation:List[Callable[[Tensor], Tensor]] = [Tensor.silu, nn.Linear(hidden_size, 2 * hidden_size, bias=True)]
+
+  def __call__(self, x:Tensor, vec:Tensor) -> Tensor:
+    shift, scale = vec.sequential(self.adaLN_modulation).chunk(2, dim=1)
+    x = (1 + scale[:, None, :]) * self.norm_final(x) + shift[:, None, :]
+    return self.linear(x)
+
+def timestep_embedding(t:Tensor, dim:int, max_period:int=10000, time_factor:float=1000.0) -> Tensor:
+  """
+  Create sinusoidal timestep embeddings.
+  :param t: a 1-D Tensor of N indices, one per batch element.
+                    These may be fractional.
+  :param dim: the dimension of the output.
+  :param max_period: controls the minimum frequency of the embeddings.
+  :return: an (N, D) Tensor of positional embeddings.
+  """
+  t = time_factor * t
+  half = dim // 2
+  freqs = Tensor.exp(-math.log(max_period) * Tensor.arange(0, stop=half, dtype=dtypes.float32) / half).to(t.device)
+
+  args = t[:, None].float() * freqs[None]
+  embedding = Tensor.cat(Tensor.cos(args), Tensor.sin(args), dim=-1)
+  if dim % 2:  embedding = Tensor.cat(*[embedding, Tensor.zeros_like(embedding[:, :1])], dim=-1)
+  if Tensor.is_floating_point(t):  embedding = embedding.cast(t.dtype)
+  return embedding
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/model.py
+class Flux:
+  """
+  Transformer model for flow matching on sequences.
+  """
+
+  def __init__(
+      self,
+      guidance_embed:bool,
+      in_channels:int = 64,
+      vec_in_dim:int = 768,
+      context_in_dim:int = 4096,
+      hidden_size:int = 3072,
+      mlp_ratio:float = 4.0,
+      num_heads:int = 24,
+      depth:int = 19,
+      depth_single_blocks:int = 38,
+      axes_dim:Optional[List[int]] = None,
+      theta:int = 10_000,
+      qkv_bias:bool = True,
+      ):
+
+    axes_dim = axes_dim or [16, 56, 56]
+    self.guidance_embed = guidance_embed
+    self.in_channels = in_channels
+    self.out_channels = self.in_channels
+    if hidden_size % num_heads != 0:
+      raise ValueError(f"Hidden size {hidden_size} must be divisible by num_heads {num_heads}")
+    pe_dim = hidden_size // num_heads
+    if sum(axes_dim) != pe_dim:
+      raise ValueError(f"Got {axes_dim} but expected positional dim {pe_dim}")
+    self.hidden_size = hidden_size
+    self.num_heads = num_heads
+    self.pe_embedder = EmbedND(dim=pe_dim, theta=theta, axes_dim=axes_dim)
+    self.img_in = nn.Linear(self.in_channels, self.hidden_size, bias=True)
+    self.time_in = MLPEmbedder(in_dim=256, hidden_dim=self.hidden_size)
+    self.vector_in = MLPEmbedder(vec_in_dim, self.hidden_size)
+    self.guidance_in:Callable[[Tensor], Tensor] = MLPEmbedder(in_dim=256, hidden_dim=self.hidden_size) if guidance_embed else tensor_identity
+    self.txt_in = nn.Linear(context_in_dim, self.hidden_size)
+
+    self.double_blocks = [DoubleStreamBlock(self.hidden_size, self.num_heads, mlp_ratio=mlp_ratio, qkv_bias=qkv_bias) for _ in range(depth)]
+    self.single_blocks = [SingleStreamBlock(self.hidden_size, self.num_heads, mlp_ratio=mlp_ratio) for _ in range(depth_single_blocks)]
+    self.final_layer = LastLayer(self.hidden_size, 1, self.out_channels)
+
+  def __call__(self, img:Tensor, img_ids:Tensor, txt:Tensor, txt_ids:Tensor, timesteps:Tensor, y:Tensor, guidance:Optional[Tensor] = None) -> Tensor:
+    if img.ndim != 3 or txt.ndim != 3:
+      raise ValueError("Input img and txt tensors must have 3 dimensions.")
+    # running on sequences img
+    img = self.img_in(img)
+    vec = self.time_in(timestep_embedding(timesteps, 256))
+    if self.guidance_embed:
+      if guidance is None:
+        raise ValueError("Didn't get guidance strength for guidance distilled model.")
+      vec = vec + self.guidance_in(timestep_embedding(guidance, 256))
+    vec = vec + self.vector_in(y)
+    txt = self.txt_in(txt)
+    ids = Tensor.cat(txt_ids, img_ids, dim=1)
+    pe = self.pe_embedder(ids)
+    for double_block in self.double_blocks:
+      img, txt = double_block(img=img, txt=txt, vec=vec, pe=pe)
+
+    img = Tensor.cat(txt, img, dim=1)
+    for single_block in self.single_blocks:
+      img = single_block(img, vec=vec, pe=pe)
+
+    img = img[:, txt.shape[1] :, ...]
+
+    return self.final_layer(img, vec)  # (N, T, patch_size ** 2 * out_channels)
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/util.py
+def load_flow_model(name:str):
+  # Loading Flux
+  print("Init model")
+  model = Flux(guidance_embed=(name != "flux-schnell"))
+  state_dict = {k.replace("scale", "weight"): v for k, v in safe_load(fetch(urls[name])).items()}
+  load_state_dict(model, state_dict)
+  return model
+
+def load_T5(max_length:int=512):
+  # max length 64, 128, 256 and 512 should work (if your sequence is short enough)
+  print("Init T5")
+  T5 = T5Embedder(max_length, fetch(urls["T5_tokenizer"]))
+  pt_1 = fetch(urls["T5_1_of_2"])
+  pt_2 = fetch(urls["T5_2_of_2"])
+  load_state_dict(T5.encoder, safe_load(pt_1) | safe_load(pt_2), strict=False)
+  return T5
+
+def load_clip():
+  print("Init Clip")
+  clip = ClipEmbedder()
+  load_state_dict(clip.transformer, safe_load(fetch(urls["clip"])))
+  return clip
+
+def load_ae() -> AutoEncoder:
+  # Loading the autoencoder
+  print("Init AE")
+  ae = AutoEncoder(0.3611, 0.1159)
+  load_state_dict(ae, safe_load(fetch(urls["ae"])))
+  return ae
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/sampling.py
+def prepare(T5:T5Embedder, clip:ClipEmbedder, img:Tensor, prompt:Union[str, List[str]]) -> Dict[str, Tensor]:
+  bs, _, h, w = img.shape
+  if bs == 1 and not isinstance(prompt, str):
+    bs = len(prompt)
+
+  img = img.rearrange("b c (h ph) (w pw) -> b (h w) (c ph pw)", ph=2, pw=2)
+  if img.shape[0] == 1 and bs > 1:
+    img = img.expand((bs, *img.shape[1:]))
+
+  img_ids = Tensor.zeros(h // 2, w // 2, 3).contiguous()
+  img_ids[..., 1] = img_ids[..., 1] + Tensor.arange(h // 2)[:, None]
+  img_ids[..., 2] = img_ids[..., 2] + Tensor.arange(w // 2)[None, :]
+  img_ids = img_ids.rearrange("h w c -> 1 (h w) c")
+  img_ids = img_ids.expand((bs, *img_ids.shape[1:]))
+
+  if isinstance(prompt, str):
+    prompt = [prompt]
+  txt = T5(prompt).realize()
+  if txt.shape[0] == 1 and bs > 1:
+    txt = txt.expand((bs, *txt.shape[1:]))
+  txt_ids = Tensor.zeros(bs, txt.shape[1], 3)
+
+  vec = clip(prompt).realize()
+  if vec.shape[0] == 1 and bs > 1:
+    vec = vec.expand((bs, *vec.shape[1:]))
+
+  return {"img": img, "img_ids": img_ids.to(img.device), "txt": txt.to(img.device), "txt_ids": txt_ids.to(img.device), "vec": vec.to(img.device)}
+
+
+def get_schedule(num_steps:int, image_seq_len:int, base_shift:float=0.5, max_shift:float=1.15, shift:bool=True) -> List[float]:
+  # extra step for zero
+  step_size = -1.0 / num_steps
+  timesteps = Tensor.arange(1, 0 + step_size, step_size)
+
+  # shifting the schedule to favor high timesteps for higher signal images
+  if shift:
+    # estimate mu based on linear estimation between two points
+    mu = 0.5 + (max_shift - base_shift) * (image_seq_len - 256) / (4096 - 256)
+    timesteps = math.exp(mu) / (math.exp(mu) + (1 / timesteps - 1))
+  return timesteps.tolist()
+
+@TinyJit
+def run(model, *args): return model(*args).realize()
+
+def denoise(model, img:Tensor, img_ids:Tensor, txt:Tensor, txt_ids:Tensor, vec:Tensor, timesteps:List[float], guidance:float=4.0) -> Tensor:
+  # this is ignored for schnell
+  guidance_vec = Tensor((guidance,), device=img.device, dtype=img.dtype).expand((img.shape[0],))
+  for t_curr, t_prev in tqdm(list(zip(timesteps[:-1], timesteps[1:])), "Denoising"):
+    t_vec = Tensor((t_curr,), device=img.device, dtype=img.dtype).expand((img.shape[0],))
+    pred = run(model, img, img_ids, txt, txt_ids, t_vec, vec, guidance_vec)
+    img = img + (t_prev - t_curr) * pred
+
+  return img
+
+def unpack(x:Tensor, height:int, width:int) -> Tensor:
+  return x.rearrange("b (h w) (c ph pw) -> b c (h ph) (w pw)", h=math.ceil(height / 16), w=math.ceil(width / 16), ph=2, pw=2)
+
+# https://github.com/black-forest-labs/flux/blob/main/src/flux/cli.py
+if __name__ == "__main__":
+  default_prompt = "bananas and a can of coke"
+  parser = argparse.ArgumentParser(description="Run Flux.1", formatter_class=argparse.ArgumentDefaultsHelpFormatter)
+
+  parser.add_argument("--name",       type=str,   default="flux-schnell", help="Name of the model to load")
+  parser.add_argument("--width",      type=int,   default=512,            help="width of the sample in pixels (should be a multiple of 16)")
+  parser.add_argument("--height",     type=int,   default=512,            help="height of the sample in pixels (should be a multiple of 16)")
+  parser.add_argument("--seed",       type=int,   default=None,           help="Set a seed for sampling")
+  parser.add_argument("--prompt",     type=str,   default=default_prompt, help="Prompt used for sampling")
+  parser.add_argument('--out',        type=str,   default=Path(tempfile.gettempdir()) / "rendered.png", help="Output filename")
+  parser.add_argument("--num_steps",  type=int,   default=None,           help="number of sampling steps (default 4 for schnell, 50 for guidance distilled)") #noqa:E501
+  parser.add_argument("--guidance",   type=float, default=3.5,            help="guidance value used for guidance distillation")
+  parser.add_argument("--output_dir", type=str,   default="output",       help="output directory")
+  args = parser.parse_args()
+
+  if args.name not in ["flux-schnell", "flux-dev"]:
+    raise ValueError(f"Got unknown model name: {args.name}, chose from flux-schnell and flux-dev")
+
+  if args.num_steps is None:
+    args.num_steps = 4 if args.name == "flux-schnell" else 50
+
+  # allow for packing and conversion to latent space
+  height = 16 * (args.height // 16)
+  width = 16 * (args.width // 16)
+
+  if args.seed is None: args.seed = Tensor._seed
+  else: Tensor.manual_seed(args.seed)
+
+  print(f"Generating with seed {args.seed}:\n{args.prompt}")
+  t0 = time.perf_counter()
+
+  # prepare input noise
+  x = Tensor.randn(1, 16, 2 * math.ceil(height / 16), 2 * math.ceil(width / 16), dtype="bfloat16")
+
+  # load text embedders
+  T5 = load_T5(max_length=256 if args.name == "flux-schnell" else 512)
+  clip = load_clip()
+
+  # embed text to get inputs for model
+  inp = prepare(T5, clip, x, prompt=args.prompt)
+  timesteps = get_schedule(args.num_steps, inp["img"].shape[1], shift=(args.name != "flux-schnell"))
+
+  # done with text embedders
+  del T5, clip
+
+  # load model
+  model = load_flow_model(args.name)
+
+  # denoise initial noise
+  x = denoise(model, **inp, timesteps=timesteps, guidance=args.guidance)
+
+  # done with model
+  del model, run
+
+  # load autoencoder
+  ae = load_ae()
+
+  # decode latents to pixel space
+  x = unpack(x.float(), height, width)
+  x = ae.decode(x).realize()
+
+  t1 = time.perf_counter()
+  print(f"Done in {t1 - t0:.1f}s. Saving {args.out}")
+
+  # bring into PIL format and save
+  x = x.clamp(-1, 1)
+  x = x[0].rearrange("c h w -> h w c")
+  x = (127.5 * (x + 1.0)).cast("uint8")
+
+  img = Image.fromarray(x.numpy())
+
+  img.save(args.out)
+
+  # validation!
+  if args.prompt == default_prompt and args.name=="flux-schnell" and args.seed == 0 and args.width == args.height == 512:
+    ref_image = Tensor(np.array(Image.open("examples/flux1_seed0.png")))
+    distance = (((x.cast(dtypes.float) - ref_image.cast(dtypes.float)) / ref_image.max())**2).mean().item()
+    assert distance < 4e-3, colored(f"validation failed with {distance=}", "red")
+    print(colored(f"output validated with {distance=}", "green"))
@@ -4,6 +4,7 @@ import argparse
 import numpy as np
 import tiktoken
 from tinygrad import Tensor, TinyJit, Device, GlobalCounters, Variable
+from tinygrad.ops import UOp
 from tinygrad.helpers import Timing, DEBUG, JIT, getenv, fetch, colored, trange
 from tinygrad.nn import Embedding, Linear, LayerNorm
 from tinygrad.nn.state import torch_load, load_state_dict, get_state_dict
@@ -75,9 +76,9 @@ class Transformer:
    self.lm_head = Linear(dim, vocab_size, bias=False)
    self.forward_jit = TinyJit(self.forward)

-  def forward(self, tokens:Union[Tensor,Variable], start_pos:Variable, temperature:float=0.0):
+  def forward(self, tokens:Union[Tensor,UOp], start_pos:Variable, temperature:float=0.0):
    if not hasattr(self, 'allpos'): self.allpos = Tensor.arange(0, MAX_CONTEXT).reshape(1, -1).realize()
-    if isinstance(tokens, Variable):
+    if isinstance(tokens, UOp):
      seqlen = 1
      tok_emb = self.wte.weight.shrink(((tokens, tokens+1), None))
    else:
@@ -107,8 +108,8 @@ class Transformer:
      ret = (logits / temperature).softmax().multinomial()
    return ret.flatten().realize()

-  def __call__(self, tokens:Tensor, start_pos:Variable, temperature:float=0.0) -> Tensor:
-    forward = (self.forward_jit if JIT and (isinstance(tokens, Variable) or tokens.shape[1] == 1) else self.forward)
+  def __call__(self, tokens:Union[Tensor,UOp], start_pos:Variable, temperature:float=0.0) -> Tensor:
+    forward = (self.forward_jit if JIT and (isinstance(tokens, UOp) or tokens.shape[1] == 1) else self.forward)
    return forward(tokens, start_pos, temperature)

 VOCAB_SIZE = 50257
@@ -1,14 +1,14 @@
-from typing import List
+from typing import List, Tuple
 from extra.models.resnet import ResNet50
 from extra.mcts_search import mcts_search
 from examples.mlperf.helpers import get_mlperf_bert_model
 from tinygrad import Tensor, Device, dtypes, nn
 from tinygrad.codegen.kernel import Kernel
+from tinygrad.ops import UOps
 from tinygrad.device import Compiled
 from tinygrad.engine.schedule import create_schedule
 from tinygrad.engine.search import time_linearizer, beam_search, bufs_from_lin
-from tinygrad.helpers import DEBUG, ansilen, getenv, colored
-from tinygrad.ops import MetaOps
+from tinygrad.helpers import DEBUG, ansilen, getenv, colored, TRACEMETA
 from tinygrad.shape.symbolic import sym_infer

 def get_sched_resnet():
@@ -17,7 +17,6 @@ def get_sched_resnet():
  BS = getenv("BS", 64)

  # run model twice to get only what changes, these are the kernels of the model
-  seen = set()
  for _ in range(2):
    out = mdl(Tensor.empty(BS, 3, 224, 224))
    targets = [out.lazydata]
@@ -25,7 +24,7 @@ def get_sched_resnet():
      optim.zero_grad()
      out.sparse_categorical_crossentropy(Tensor.empty(BS, dtype=dtypes.int)).backward()
      targets += [x.lazydata for x in optim.schedule_step()]
-    sched = create_schedule(targets, seen)
+    sched = create_schedule(targets)
    print(f"schedule length {len(sched)}")
  return sched

@@ -34,7 +33,7 @@ def get_sched_bert():
  optim = nn.optim.LAMB(nn.state.get_parameters(mdl))

  # fake data
-  BS = getenv("BS", 2)
+  BS = getenv("BS", 9)
  input_ids = Tensor.empty((BS, 512), dtype=dtypes.float32)
  segment_ids = Tensor.empty((BS, 512), dtype=dtypes.float32)
  attention_mask = Tensor.empty((BS, 512), dtype=dtypes.default_float)
@@ -54,7 +53,7 @@ def get_sched_bert():
      # ignore grad norm and loss scaler for now
      loss.backward()
      targets += [x.lazydata for x in optim.schedule_step()]
-    sched = create_schedule(targets, seen)
+    sched = create_schedule(targets)
    print(f"schedule length {len(sched)}")
  return sched

@@ -68,7 +67,7 @@ if __name__ == "__main__":
  print(f"optimizing for {Device.DEFAULT}")

  sched = globals()[f"get_sched_{getenv('MODEL', 'resnet')}"]()
-  sched = [x for x in sched if x.ast.op is MetaOps.KERNEL]
+  sched = [x for x in sched if x.ast.op is UOps.SINK]

  # focus on one kernel
  if getenv("KERNEL", -1) >= 0: sched = sched[getenv("KERNEL", -1):getenv("KERNEL", -1)+1]
@@ -83,7 +82,7 @@ if __name__ == "__main__":
    rawbufs = bufs_from_lin(Kernel(si.ast))

    # "linearize" the op into uops in different ways
-    lins:List[Kernel] = []
+    lins: List[Tuple[Kernel, str]] = []

    # always try hand coded opt
    lin = Kernel(si.ast, opts=device.renderer)
@@ -109,10 +108,10 @@ if __name__ == "__main__":

    # benchmark the programs
    choices = []
-    for (lin, nm) in lins:
+    for lin, nm in lins:
      tm = time_linearizer(lin, rawbufs, allow_test_size=False, cnt=10, disable_cache=True)
      ops = (prg:=lin.to_program()).op_estimate
-      gflops = sym_infer(ops, {k:k.min for k in lin.ast.vars()})*1e-9/tm
+      gflops = sym_infer(ops, {k:k.min for k in lin.ast.variables()})*1e-9/tm
      choices.append((tm, gflops, lin, prg, nm))

    sorted_choices = sorted(choices, key=lambda x: x[0])
@@ -129,7 +128,7 @@ if __name__ == "__main__":
    running_gflops += gflops * tm
    if (key := str([str(m) for m in si.metadata] if si.metadata is not None else None)) not in usage: usage[key] = (0, 0)
    usage[key] = (usage[key][0] + tm, usage[key][1] + 1)
-    print(f"*** {total_tm*1000:7.2f} ms : kernel {i:2d} {lin.name+' '*(37-ansilen(lin.name))} {str(prg.global_size):18s} {str(prg.local_size):12s} takes {tm*1000:7.2f} ms, {gflops:6.0f} GFLOPS {[str(m) for m in si.metadata] if si.metadata is not None else ''}")
+    print(f"*** {total_tm*1000:7.2f} ms : kernel {i:2d} {lin.name+' '*(37-ansilen(lin.name))} {str(prg.global_size):18s} {str(prg.local_size):12s} takes {tm*1000:7.2f} ms, {gflops:6.0f} GFLOPS {[repr(m) if TRACEMETA >= 2 else str(m) for m in si.metadata] if si.metadata is not None else ''}")
  print(f"******* total {total_tm*1000:.2f} ms, {running_gflops/total_tm:6.0f} GFLOPS")
  print("usage:")
  for k in sorted(usage, key=lambda x: -usage[x][0])[:10]:
@@ -6,7 +6,6 @@
 import random, time
 import numpy as np
 from typing import Optional
-from extra.datasets import fetch_cifar, cifar_mean, cifar_std
 from extra.lr_scheduler import OneCycleLR
 from tinygrad import nn, dtypes, Tensor, Device, GlobalCounters, TinyJit
 from tinygrad.nn.state import get_state_dict, get_parameters
@@ -14,6 +13,9 @@ from tinygrad.nn import optim
 from tinygrad.helpers import Context, BEAM, WINO, getenv, colored, prod
 from tinygrad.multi import MultiLazyBuffer

+cifar_mean = [0.4913997551666284, 0.48215855929893703, 0.4465309133731618]
+cifar_std = [0.24703225141799082, 0.24348516474564, 0.26158783926049628]
+
 BS, STEPS = getenv("BS", 512), getenv("STEPS", 1000)
 EVAL_BS = getenv("EVAL_BS", BS)
 GPUS = [f'{Device.DEFAULT}:{i}' for i in range(getenv("GPUS", 1))]
@@ -252,7 +254,7 @@ def train_cifar():
      if not is_train: break

  transform = [
-    lambda x: x / 255.0,
+    lambda x: x.float() / 255.0,
    lambda x: x.reshape((-1,3,32,32)) - Tensor(cifar_mean, device=x.device, dtype=x.dtype).reshape((1,3,1,1)),
    lambda x: x / Tensor(cifar_std, device=x.device, dtype=x.dtype).reshape((1,3,1,1)),
  ]
@@ -277,10 +279,7 @@ def train_cifar():

  set_seed(getenv('SEED', hyp['seed']))

-  X_train, Y_train, X_test, Y_test = fetch_cifar()
-  # load data and label into GPU and convert to dtype accordingly
-  X_train, X_test = X_train.to(device=Device.DEFAULT).float(), X_test.to(device=Device.DEFAULT).float()
-  Y_train, Y_test = Y_train.to(device=Device.DEFAULT), Y_test.to(device=Device.DEFAULT)
+  X_train, Y_train, X_test, Y_test = nn.datasets.cifar()
  # one-hot encode labels
  Y_train, Y_test = Y_train.one_hot(10), Y_test.one_hot(10)
  # preprocess data
@@ -142,8 +142,7 @@ def build_transformer(model_path: Path, model_size="8B", quantize=None, device=N
  if quantize == "int8": linear = Int8Linear
  elif quantize == "nf4": linear = NF4Linear(64)
  else: linear = nn.Linear
-  with Context(THREEFRY=0):
-    model = Transformer(**MODEL_PARAMS[model_size]["args"], linear=linear, max_context=8192, jit=True)
+  model = Transformer(**MODEL_PARAMS[model_size]["args"], linear=linear, max_context=8192, jit=True)

  # load weights
  if model_path.is_dir():
@@ -7,7 +7,7 @@ from train_gpt2 import GPT, GPTConfig
 from tinygrad.helpers import dedup, to_function_name, flatten, getenv, GRAPH, GlobalCounters, ansilen, to_function_name
 from tinygrad.engine.schedule import create_schedule
 from tinygrad.engine.realize import get_kernel, memory_planner, run_schedule
-from tinygrad.ops import BufferOps, MetaOps
+from tinygrad.ops import MetaOps, UOps

 TIMING = getenv("TIMING")

@@ -16,8 +16,7 @@ if __name__ == "__main__":
  #model.load_pretrained()
  for p in nn.state.get_parameters(model): p.replace(Tensor.empty(p.shape, dtype=p.dtype)) # fake load pretrained

-  seen = set()
-  #early_sched = create_schedule([x.lazydata for x in nn.state.get_parameters(model)], seen)
+  #early_sched = create_schedule([x.lazydata for x in nn.state.get_parameters(model)])
  #print(f"built model {len(early_sched)}")

  #B, T = Variable("B", 1, 128).bind(4), 64 #Variable("T", 1, 1024).bind(64)
@@ -38,12 +37,11 @@ if __name__ == "__main__":
      tensors = optimizer.schedule_step()
    else:
      tensors = []
-    sched = create_schedule([loss.lazydata] + [x.lazydata for x in tensors], seen)
+    sched = create_schedule([loss.lazydata] + [x.lazydata for x in tensors])
    print(f"calls {i}:", len(sched))
    #run_schedule(sched[:])
-  del seen  # free the LazyBuffers
  sched = memory_planner(sched)
-  ast_dedup = dedup([si.ast for si in sched if si.ast[0].op is BufferOps.STORE])
+  ast_dedup = dedup([si.ast for si in sched if si.ast.op is UOps.SINK])
  srcs = {}
  for ast in ast_dedup:
    k = get_kernel(Device["CLANG"].renderer, ast)
@@ -84,8 +82,8 @@ if __name__ == "__main__":
  for i,si in enumerate(sched):
    bufs = [(named_buffers.get(b, f"b{numbered_bufs[b]}"), b) for b in si.bufs]
    all_bufs += bufs
-    if si.ast[0].op is not BufferOps.STORE:
-      print(f"// {si.ast[0].op}", bufs)
+    if si.ast.op is not UOps.SINK:
+      print(f"// {si.ast.op}", bufs)
    else:
      print(f"{srcs[si.ast][0]}({', '.join([x[0] for x in bufs])})")
      main.append(f"  {to_function_name(srcs[si.ast][0])}({', '.join([x[0] for x in bufs])});")
@@ -8,6 +8,7 @@ from dataclasses import dataclass
 class GPTConfig:
  block_size: int = 1024
  vocab_size: int = 50257
+  padded_vocab_size: int = 50304
  n_layer: int = 12
  n_head: int = 12
  n_embd: int = 768
@@ -68,19 +69,21 @@ class GPT:
  def __init__(self, config:GPTConfig):
    self.config = config

-    self.wte = nn.Embedding(config.vocab_size, config.n_embd)
+    self.wte = nn.Embedding(config.padded_vocab_size, config.n_embd)
    self.wpe = nn.Embedding(config.block_size, config.n_embd)
    self.h = [Block(config) for _ in range(config.n_layer)]
    self.ln_f = nn.LayerNorm(config.n_embd)
-    self.lm_head = nn.Linear(config.n_embd, config.vocab_size, bias=False)
+    self.lm_head = nn.Linear(config.n_embd, config.padded_vocab_size, bias=False)
    self.wte.weight = self.lm_head.weight # https://paperswithcode.com/method/weight-tying

  def load_pretrained(self):
    weights = nn.state.torch_load(fetch(f'https://huggingface.co/gpt2/resolve/main/pytorch_model.bin'))
    transposed = ('attn.c_attn.weight', 'attn.c_proj.weight', 'mlp.c_fc.weight', 'mlp.c_proj.weight')
    for k in weights:
+      if k == "wte.weight":
+        weights[k] = weights[k].pad(((0, self.config.padded_vocab_size-self.config.vocab_size), (0,0))).to(None).contiguous()
      if k.endswith(transposed):
-        weights[k] = weights[k].to(Device.DEFAULT).T.contiguous()
+        weights[k] = weights[k].to(None).T.contiguous()
    # lm head and wte are tied
    weights['lm_head.weight'] = weights['wte.weight']
    nn.state.load_state_dict(self, weights)
@@ -105,10 +108,10 @@ class GPT:
    x = self.ln_f(x.sequential(self.h))

    if targets is not None:
-      logits = self.lm_head(x)
+      logits = self.lm_head(x)[:, :, :self.config.vocab_size]
      loss = logits.sparse_categorical_crossentropy(targets)
    else:
-      logits = self.lm_head(x[:, [-1], :])
+      logits = self.lm_head(x[:, [-1], :])[:, :, :self.config.vocab_size]
      loss = None

    return logits, loss
@@ -120,6 +123,7 @@ if __name__ == "__main__":
  parser.add_argument("--num_iterations", type=int, default=10, help="number of iterations to run")
  parser.add_argument("--batch_size", type=int, default=4, help="batch size")
  parser.add_argument("--sequence_length", type=int, default=64, help="sequence length")
+  parser.add_argument("--skip_test", action="store_true", help="skip test")
  args = parser.parse_args()
  B, T = args.batch_size, args.sequence_length
  assert 1 <= T <= 1024
@@ -135,10 +139,7 @@ if __name__ == "__main__":
  # load the tokens
  # prefer to use tiny_shakespeare if it's available, otherwise use tiny_stories
  # we're using val instead of train split just because it is smaller/faster
-  shake_tokens_bin = "data/tiny_shakespeare_val.bin"
-  story_tokens_bin = "data/TinyStories_val.bin"
-  assert os.path.isfile(shake_tokens_bin) or os.path.isfile(story_tokens_bin), "you must run prepro on some dataset"
-  tokens_bin = shake_tokens_bin if os.path.isfile(shake_tokens_bin) else story_tokens_bin
+  tokens_bin = fetch("https://huggingface.co/datasets/karpathy/llmc-starter-pack/resolve/main/tiny_shakespeare_val.bin")
  assert os.path.isfile(tokens_bin)
  print(f"loading cached tokens in {tokens_bin}")
  with open(tokens_bin, "rb") as f:
@@ -169,8 +170,7 @@ if __name__ == "__main__":
    _, loss = model(x, y)
    optimizer.zero_grad()
    loss.backward()
-    optimizer.step()
-    return loss
+    return loss.realize(*optimizer.schedule_step())

  with Tensor.train():
    for i in range(args.num_iterations):
@@ -179,14 +179,15 @@ if __name__ == "__main__":
      loss = step(x.contiguous(), y.contiguous())
      Device[Device.DEFAULT].synchronize()
      t1 = time.time()
-      print(f"iteration {i}, loss: {loss.item()}, time: {(t1-t0)*1000:.3f}ms")
+      print(f"iteration {i}, loss: {loss.item():.6f}, time: {(t1-t0)*1000:.3f}ms, {int(B*T/(t1-t0))} tok/s")

-  start = "<|endoftext|>"
-  start_ids = encode(start)
-  x = (Tensor(start_ids)[None, ...])
-  max_new_tokens = 16
-  temperature = 1.0
-  top_k = 40
-  y = model.generate(x, max_new_tokens, temperature=temperature, top_k=top_k)
-  print(decode(y[0].tolist()))
+  if not args.skip_test:
+    start = "<|endoftext|>"
+    start_ids = encode(start)
+    x = (Tensor(start_ids)[None, ...])
+    max_new_tokens = 16
+    temperature = 1.0
+    top_k = 40
+    y = model.generate(x, max_new_tokens, temperature=temperature, top_k=top_k)
+    print(decode(y[0].tolist()))

@@ -94,11 +94,11 @@ def eval_retinanet():
    x /= input_std
    return x

-  from extra.datasets.openimages import openimages, iterate
+  from extra.datasets.openimages import download_dataset, iterate, BASEDIR
  from pycocotools.coco import COCO
  from pycocotools.cocoeval import COCOeval
  from contextlib import redirect_stdout
-  coco = COCO(openimages('validation'))
+  coco = COCO(download_dataset(base_dir:=getenv("BASE_DIR", BASEDIR), 'validation'))
  coco_eval = COCOeval(coco, iouType="bbox")
  coco_evalimgs, evaluated_imgs, ncats, narea = [], [], len(coco_eval.params.catIds), len(coco_eval.params.areaRng)

@@ -107,13 +107,13 @@ def eval_retinanet():

  n, bs = 0, 8
  st = time.perf_counter()
-  for x, targets in iterate(coco, bs):
+  for x, targets in iterate(coco, base_dir, bs):
    dat = Tensor(x.astype(np.float32))
    mt = time.perf_counter()
    if dat.shape[0] == bs:
      outs = mdlrun(dat).numpy()
    else:
-      mdlrun.jit_cache = None
+      mdlrun._jit_cache = []
      outs =  mdl(input_fixup(dat)).numpy()
    et = time.perf_counter()
    predictions = mdl.postprocess_detections(outs, input_size=dat.shape[1:3], orig_image_sizes=[t["image_size"] for t in targets])
@@ -181,7 +181,7 @@ def eval_bert():
  from examples.mlperf.metrics import f1_score
  from transformers import BertTokenizer

-  tokenizer = BertTokenizer(str(Path(__file__).parents[2] / "weights/bert_vocab.txt"))
+  tokenizer = BertTokenizer(str(Path(__file__).parents[2] / "extra/weights/bert_vocab.txt"))

  c = 0
  f1 = 0.0
@@ -4,7 +4,7 @@ from tqdm import tqdm
 import multiprocessing

 from tinygrad import Device, GlobalCounters, Tensor, TinyJit, dtypes
-from tinygrad.helpers import getenv, BEAM, WINO, round_up, diskcache_clear
+from tinygrad.helpers import getenv, BEAM, WINO, round_up, diskcache_clear, FUSE_CONV_BW
 from tinygrad.nn.state import get_parameters, get_state_dict, safe_load, safe_save
 from tinygrad.nn.optim import LAMB, LARS, SGD, OptimizerGroup

@@ -28,7 +28,7 @@ def train_resnet():
  if getenv("LOGMLPERF"):
    from mlperf_logging import mllog
    import mlperf_logging.mllog.constants as mllog_constants
-    mllog.config(filename=f"result_{seed}.txt")
+    mllog.config(filename=f"result_resnet_{seed}.txt")
    mllog.config(root_dir=Path(__file__).parents[3].as_posix())  # truncate to log this. "file": "tinygrad/examples/mlperf/model_train.py"
    MLLOGGER = mllog.get_mllogger()
    if INITMLPERF:
@@ -346,8 +346,225 @@ def train_retinanet():
  pass

 def train_unet3d():
-  # TODO: Unet3d
-  pass
+  """
+  Trains the UNet3D model.
+
+  Instructions:
+  1) Run the following script from the root folder of `tinygrad`:
+  ```./examples/mlperf/scripts/setup_kits19_dataset.sh```
+
+  Optionally, `BASEDIR` can be set to download and process the dataset at a specific location:
+  ```BASEDIR=<folder_path> ./examples/mlperf/scripts/setup_kits19_dataset.sh```
+
+  2) To start training the model, run the following:
+  ```time PYTHONPATH=. WANDB=1 TRAIN_BEAM=3 FUSE_CONV_BW=1 GPUS=6 BS=6 MODEL=unet3d python3 examples/mlperf/model_train.py```
+  """
+  from examples.mlperf.losses import dice_ce_loss
+  from examples.mlperf.metrics import dice_score
+  from examples.mlperf.dataloader import batch_load_unet3d
+  from extra.models.unet3d import UNet3D
+  from extra.datasets.kits19 import iterate, get_train_files, get_val_files, sliding_window_inference, preprocess_dataset, TRAIN_PREPROCESSED_DIR, VAL_PREPROCESSED_DIR
+  from tinygrad import Context
+  from tinygrad.nn.optim import SGD
+  from math import ceil
+
+  GPUS = [f"{Device.DEFAULT}:{i}" for i in range(getenv("GPUS", 1))]
+  for x in GPUS: Device[x]
+
+  TARGET_METRIC = 0.908
+  NUM_EPOCHS = getenv("NUM_EPOCHS", 4000)
+  BS = getenv("BS", 1 * len(GPUS))
+  LR = getenv("LR", 2.0 * (BS / 28))
+  LR_WARMUP_EPOCHS = getenv("LR_WARMUP_EPOCHS", 1000)
+  LR_WARMUP_INIT_LR = getenv("LR_WARMUP_INIT_LR", 0.0001)
+  WANDB = getenv("WANDB")
+  PROJ_NAME = getenv("PROJ_NAME", "tinygrad_unet3d_mlperf")
+  SEED = getenv("SEED", -1) if getenv("SEED", -1) >= 0 else None
+  TRAIN_DATASET_SIZE, VAL_DATASET_SIZE = len(get_train_files()), len(get_val_files())
+  SAMPLES_PER_EPOCH = TRAIN_DATASET_SIZE // BS
+  START_EVAL_AT = getenv("START_EVAL_AT", ceil(1000 * TRAIN_DATASET_SIZE / (SAMPLES_PER_EPOCH * BS)))
+  EVALUATE_EVERY = getenv("EVALUATE_EVERY", ceil(20 * TRAIN_DATASET_SIZE / (SAMPLES_PER_EPOCH * BS)))
+  TRAIN_BEAM, EVAL_BEAM = getenv("TRAIN_BEAM", BEAM.value), getenv("EVAL_BEAM", BEAM.value)
+  BENCHMARK = getenv("BENCHMARK")
+  CKPT = getenv("CKPT")
+
+  config = {
+    "num_epochs": NUM_EPOCHS,
+    "batch_size": BS,
+    "learning_rate": LR,
+    "learning_rate_warmup_epochs": LR_WARMUP_EPOCHS,
+    "learning_rate_warmup_init": LR_WARMUP_INIT_LR,
+    "start_eval_at": START_EVAL_AT,
+    "evaluate_every": EVALUATE_EVERY,
+    "train_beam": TRAIN_BEAM,
+    "eval_beam": EVAL_BEAM,
+    "wino": WINO.value,
+    "fuse_conv_bw": FUSE_CONV_BW.value,
+    "gpus": GPUS,
+    "default_float": dtypes.default_float.name
+  }
+
+  if WANDB:
+    try:
+      import wandb
+    except ImportError:
+      raise "Need to install wandb to use it"
+
+  if SEED is not None:
+    config["seed"] = SEED
+    Tensor.manual_seed(SEED)
+
+  model = UNet3D()
+  params = get_parameters(model)
+
+  for p in params: p.realize().to_(GPUS)
+
+  optim = SGD(params, lr=LR, momentum=0.9, nesterov=True)
+
+  def lr_warm_up(optim, init_lr, lr, current_epoch, warmup_epochs):
+    scale = current_epoch / warmup_epochs
+    optim.lr.assign(Tensor([init_lr + (lr - init_lr) * scale], device=GPUS)).realize()
+
+  def save_checkpoint(state_dict, fn):
+    if not os.path.exists("./ckpts"): os.mkdir("./ckpts")
+    print(f"saving checkpoint to {fn}")
+    safe_save(state_dict, fn)
+
+  def data_get(it):
+    x, y, cookie = next(it)
+    return x.shard(GPUS, axis=0).realize(), y.shard(GPUS, axis=0), cookie
+
+  @TinyJit
+  @Tensor.train()
+  def train_step(model, x, y):
+    optim.zero_grad()
+
+    y_hat = model(x)
+    loss = dice_ce_loss(y_hat, y)
+
+    loss.backward()
+    optim.step()
+    return loss.realize()
+  
+  @Tensor.train(mode=False)
+  @Tensor.test()
+  def eval_step(model, x, y):
+    y_hat, y = sliding_window_inference(model, x, y, gpus=GPUS)
+    y_hat, y = Tensor(y_hat), Tensor(y, requires_grad=False)
+    loss = dice_ce_loss(y_hat, y)
+    score = dice_score(y_hat, y)
+    return loss.realize(), score.realize()
+  
+  if WANDB: wandb.init(config=config, project=PROJ_NAME)
+
+  step_times, start_epoch = [], 1
+  is_successful, diverged = False, False
+  start_eval_at, evaluate_every = 1 if BENCHMARK else START_EVAL_AT, 1 if BENCHMARK else EVALUATE_EVERY
+  next_eval_at = start_eval_at
+
+  print(f"Training on {GPUS}")
+  
+  if BENCHMARK: print("Benchmarking UNet3D")
+  else: print(f"Start evaluation at epoch {start_eval_at} and every {evaluate_every} epoch(s) afterwards")
+
+  if not TRAIN_PREPROCESSED_DIR.exists(): preprocess_dataset(get_train_files(), TRAIN_PREPROCESSED_DIR, False)
+  if not VAL_PREPROCESSED_DIR.exists(): preprocess_dataset(get_val_files(), VAL_PREPROCESSED_DIR, True)
+
+  for epoch in range(1, NUM_EPOCHS + 1):
+    with Context(BEAM=TRAIN_BEAM):
+      if epoch <= LR_WARMUP_EPOCHS and LR_WARMUP_EPOCHS > 0:
+        lr_warm_up(optim, LR_WARMUP_INIT_LR, LR, epoch, LR_WARMUP_EPOCHS)
+
+      train_dataloader = batch_load_unet3d(TRAIN_PREPROCESSED_DIR, batch_size=BS, val=False, shuffle=True, seed=SEED)
+      it = iter(tqdm(train_dataloader, total=SAMPLES_PER_EPOCH, desc=f"epoch {epoch}", disable=BENCHMARK))
+      i, proc = 0, data_get(it)
+
+      prev_cookies = []
+      st = time.perf_counter()
+
+      while proc is not None:
+        GlobalCounters.reset()
+
+        loss, proc = train_step(model, proc[0], proc[1]), proc[2]
+
+        pt = time.perf_counter()
+
+        if len(prev_cookies) == getenv("STORE_COOKIES", 1): prev_cookies = []  # free previous cookies after gpu work has been enqueued
+        try:
+          next_proc = data_get(it)
+        except StopIteration:
+          next_proc = None
+
+        dt = time.perf_counter()
+
+        device_str = loss.device if isinstance(loss.device, str) else f"{loss.device[0]} * {len(loss.device)}"
+        loss = loss.numpy().item()
+
+        cl = time.perf_counter()
+
+        if BENCHMARK: step_times.append(cl - st)
+
+        tqdm.write(
+          f"{i:5} {((cl - st)) * 1000.0:7.2f} ms run, {(pt - st) * 1000.0:7.2f} ms python, {(dt - pt) * 1000.0:6.2f} ms fetch data, "
+          f"{(cl - dt) * 1000.0:7.2f} ms {device_str}, {loss:5.2f} loss, {optim.lr.numpy()[0]:.6f} LR, "
+          f"{GlobalCounters.mem_used / 1e9:.2f} GB used, {GlobalCounters.global_ops * 1e-9 / (cl - st):9.2f} GFLOPS"
+        )
+
+        if WANDB:
+          wandb.log({"lr": optim.lr.numpy(), "train/loss": loss, "train/step_time": cl - st, "train/python_time": pt - st, "train/data_time": dt - pt,
+                     "train/cl_time": cl - dt, "train/GFLOPS": GlobalCounters.global_ops * 1e-9 / (cl - st), "epoch": epoch + (i + 1) / SAMPLES_PER_EPOCH})
+
+        st = cl
+        prev_cookies.append(proc)
+        proc, next_proc = next_proc, None  # return old cookie
+        i += 1
+
+        if i == BENCHMARK:
+          median_step_time = sorted(step_times)[(BENCHMARK + 1) // 2]  # in seconds
+          estimated_total_minutes = int(median_step_time * SAMPLES_PER_EPOCH * NUM_EPOCHS / 60)
+          print(f"Estimated training time: {estimated_total_minutes // 60}h{estimated_total_minutes % 60}m")
+          if (TRAIN_BEAM or EVAL_BEAM) and epoch == start_epoch: break
+          return
+
+    with Context(BEAM=EVAL_BEAM):
+      if epoch == next_eval_at:
+        next_eval_at += evaluate_every
+        eval_loss = []
+        scores = []
+
+        for x, y in tqdm(iterate(get_val_files(), preprocessed_dir=VAL_PREPROCESSED_DIR), total=VAL_DATASET_SIZE):
+          eval_loss_value, score = eval_step(model, x, y)
+          eval_loss.append(eval_loss_value)
+          scores.append(score)
+
+        scores = Tensor.mean(Tensor.stack(*scores, dim=0), axis=0).numpy()
+        eval_loss = Tensor.mean(Tensor.stack(*eval_loss, dim=0), axis=0).numpy()
+
+        l1_dice, l2_dice = scores[0][-2], scores[0][-1]
+        mean_dice = (l2_dice + l1_dice) / 2
+
+        tqdm.write(f"{l1_dice} L1 dice, {l2_dice} L2 dice, {mean_dice:.3f} mean_dice, {eval_loss:5.2f} eval_loss")
+
+        if WANDB:
+          wandb.log({"eval/loss": eval_loss, "eval/mean_dice": mean_dice, "epoch": epoch})
+
+        if mean_dice >= TARGET_METRIC:
+          is_successful = True
+          save_checkpoint(get_state_dict(model), f"./ckpts/unet3d.safe")
+        elif mean_dice < 1e-6:
+          print("Model diverging. Aborting.")
+          diverged = True
+
+    if not is_successful and CKPT:
+      if WANDB and wandb.run is not None:
+        fn = f"./ckpts/{time.strftime('%Y%m%d_%H%M%S')}_{wandb.run.id}_e{epoch}.safe"
+      else:
+        fn = f"./ckpts/{time.strftime('%Y%m%d_%H%M%S')}_e{epoch}.safe"
+
+      save_checkpoint(get_state_dict(model), fn)
+
+    if is_successful or diverged:
+      break

 def train_rnnt():
  # TODO: RNN-T
@@ -404,7 +621,7 @@ def train_bert():
    from mlperf_logging import mllog
    import mlperf_logging.mllog.constants as mllog_constants

-    mllog.config(filename="bert.log")
+    mllog.config(filename=f"result_bert_{seed}.log")
    mllog.config(root_dir=Path(__file__).parents[3].as_posix())
    MLLOGGER = mllog.get_mllogger()
    MLLOGGER.logger.propagate = False
@@ -432,7 +649,7 @@ def train_bert():
  EVAL_BS            = config["EVAL_BS"]                = getenv("EVAL_BS", 1 * len(GPUS))
  max_lr             = config["OPT_BASE_LEARNING_RATE"] = getenv("OPT_BASE_LEARNING_RATE", 0.0001 * math.sqrt(BS/66))

-  train_steps        = config["TRAIN_STEPS"]            = getenv("TRAIN_STEPS", 3000000 // BS)
+  train_steps        = config["TRAIN_STEPS"]            = getenv("TRAIN_STEPS", 3630000 // BS)
  warmup_steps       = config["NUM_WARMUP_STEPS"]       = getenv("NUM_WARMUP_STEPS", 1)
  max_eval_steps     = config["MAX_EVAL_STEPS"]         = getenv("MAX_EVAL_STEPS", (10000 + EVAL_BS - 1) // EVAL_BS) # EVAL_BS * MAX_EVAL_STEPS >= 10000
  eval_step_freq     = config["EVAL_STEP_FREQ"]         = getenv("EVAL_STEP_FREQ", int((math.floor(0.05 * (230.23 * BS + 3000000) / 25000) * 25000) / BS)) # Round down
@@ -441,7 +658,7 @@ def train_bert():
  save_ckpt_dir      = config["SAVE_CKPT_DIR"]          = getenv("SAVE_CKPT_DIR", "./ckpts")
  init_ckpt          = config["INIT_CKPT_DIR"]          = getenv("INIT_CKPT_DIR", BASEDIR)

-  loss_scaler        = config["LOSS_SCALER"]            = getenv("LOSS_SCALER", 2.0**13 if dtypes.default_float == dtypes.float16 else 1.0)
+  loss_scaler        = config["LOSS_SCALER"]            = getenv("LOSS_SCALER", 2.0**10 if dtypes.default_float == dtypes.float16 else 1.0)
  decay              = config["DECAY"]                  = getenv("DECAY", 0.01)
  epsilon            = config["EPSILON"]                = getenv("EPSILON", 1e-6)
  poly_power         = config["POLY_POWER"]             = getenv("POLY_POWER", 1.0)
@@ -455,14 +672,23 @@ def train_bert():

  Tensor.manual_seed(seed)  # seed for weight initialization

-  model = get_mlperf_bert_model(init_ckpt if not INITMLPERF else None)
+  assert 10000 <= (EVAL_BS * max_eval_steps), "Evaluation batchsize * max_eval_steps must greater or equal 10000 to iterate over full eval dataset"
+
+  # ** init wandb **
+  WANDB = getenv("WANDB")
+  if WANDB:
+    import wandb
+    wandb_args = {"id": wandb_id, "resume": "must"} if (wandb_id := getenv("WANDB_RESUME", "")) else {}
+    wandb.init(config=config, **wandb_args, project="MLPerf-BERT")
+
+  # ** init model **
+
+  model = get_mlperf_bert_model(init_ckpt if RUNMLPERF else None)
  
  for _, x in get_state_dict(model).items():
    x.realize().to_(GPUS)
  parameters = get_parameters(model)

-  assert 10000 <= (EVAL_BS * max_eval_steps), "Evaluation batchsize * max_eval_steps must greater or equal 10000 to iterate over full eval dataset"
-
  # ** Log run config **
  for key, value in config.items(): print(f'HParam: "{key}": {value}')

@@ -510,14 +736,8 @@ def train_bert():
    start_step = int(scheduler_wd.epoch_counter.numpy().item())
    print(f"resuming from {ckpt} at step {start_step}")

-  # ** init wandb **
-  WANDB = getenv("WANDB")
-  if WANDB:
-    import wandb
-    wandb_args = {"id": wandb_id, "resume": "must"} if (wandb_id := getenv("WANDB_RESUME", "")) else {}
-    wandb.init(config=config, **wandb_args, project="MLPerf-BERT")
-
-  if not INITMLPERF:
+  if RUNMLPERF:
+    # only load real data with RUNMLPERF
    eval_it = iter(batch_load_val_bert(EVAL_BS))
    train_it = iter(tqdm(batch_load_train_bert(BS), total=train_steps, disable=BENCHMARK))
    for _ in range(start_step): next(train_it) # Fast forward
@@ -526,10 +746,14 @@ def train_bert():
  step_times = []
  # ** train loop **
  wc_start = time.perf_counter()
-  if INITMLPERF:
-    i, train_data = start_step, get_fake_data_bert(GPUS, BS)
-  else:
+  if RUNMLPERF:
+    # only load real data with RUNMLPERF
    i, train_data = start_step, get_data_bert(GPUS, train_it)
+    if MLLOGGER:
+      MLLOGGER.start(key=mllog_constants.EPOCH_START, value=i*BS, metadata={"epoch_num": i*BS})
+  else:
+    i, train_data = start_step, get_fake_data_bert(GPUS, BS)
+
  while train_data is not None and i < train_steps and not achieved:
    Tensor.training = True
    BEAM.value = TRAIN_BEAM
@@ -542,10 +766,10 @@ def train_bert():
    pt = time.perf_counter()

    try:
-      if INITMLPERF:
-        next_data = get_fake_data_bert(GPUS, BS)
-      else:
+      if RUNMLPERF:
        next_data = get_data_bert(GPUS, train_it)
+      else:
+        next_data = get_fake_data_bert(GPUS, BS)
    except StopIteration:
      next_data = None

@@ -564,7 +788,7 @@ def train_bert():
    if WANDB:
      wandb.log({"lr": optimizer_wd.lr.numpy(), "train/loss": loss, "train/step_time": cl - st,
                  "train/python_time": pt - st, "train/data_time": dt - pt, "train/cl_time": cl - dt,
-                  "train/GFLOPS": GlobalCounters.global_ops * 1e-9 / (cl - st)})
+                  "train/GFLOPS": GlobalCounters.global_ops * 1e-9 / (cl - st), "epoch": (i+1)*BS})

    train_data, next_data = next_data, None
    i += 1
@@ -579,8 +803,8 @@ def train_bert():
    # ** eval loop **
    if i % eval_step_freq == 0 or (BENCHMARK and i == BENCHMARK):
      if MLLOGGER and RUNMLPERF:
-        MLLOGGER.start(key=mllog_constants.EVAL_START, value=None, metadata={"epoch_num": 1, "epoch_count": 1, "step_num": i})
-      train_step_bert.reset()
+        MLLOGGER.start(key=mllog_constants.EVAL_START, value=None, metadata={"epoch_num": i*BS, "step_num": i})
+      if getenv("RESET_STEP", 1): train_step_bert.reset()
      eval_lm_losses = []
      eval_clsf_losses = []
      eval_lm_accs = []
@@ -590,10 +814,10 @@ def train_bert():
      BEAM.value = EVAL_BEAM

      for j in tqdm(range(max_eval_steps), desc="Evaluating", total=max_eval_steps, disable=BENCHMARK):
-        if INITMLPERF:
-          eval_data = get_fake_data_bert(GPUS, EVAL_BS)
-        else:
+        if RUNMLPERF:
          eval_data = get_data_bert(GPUS, eval_it)
+        else:
+          eval_data = get_fake_data_bert(GPUS, EVAL_BS)
        GlobalCounters.reset()
        st = time.time()

@@ -618,7 +842,8 @@ def train_bert():
            MLLOGGER.event(key=mllog_constants.INIT_STOP, value=None)
          return

-      eval_step_bert.reset()
+      if getenv("RESET_STEP", 1): eval_step_bert.reset()
+      del eval_data, eval_result
      avg_lm_loss = sum(eval_lm_losses) / len(eval_lm_losses)
      avg_clsf_loss = sum(eval_clsf_losses) / len(eval_clsf_losses)
      avg_lm_acc = sum(eval_lm_accs) / len(eval_lm_accs)
@@ -633,16 +858,17 @@ def train_bert():
                    "eval/clsf_accuracy": avg_clsf_acc, "eval/forward_time": avg_fw_time})

      if MLLOGGER and RUNMLPERF:
-        MLLOGGER.end(key=mllog_constants.EVAL_STOP, value=i, metadata={"epoch_count": 1, "step_num": i, "samples_count": config["EVAL_BS"] * config["MAX_EVAL_STEPS"]})
-        MLLOGGER.event(key=mllog_constants.EVAL_ACCURACY, value=avg_lm_acc, metadata={"epoch_num": 1, "masked_lm_accuracy": avg_lm_acc})
+        MLLOGGER.end(key=mllog_constants.EVAL_STOP, value=i*BS, metadata={"epoch_count": i*BS, "step_num": i, "samples_count": config["EVAL_BS"] * config["MAX_EVAL_STEPS"]})
+        MLLOGGER.event(key=mllog_constants.EVAL_ACCURACY, value=avg_lm_acc, metadata={"epoch_num": i*BS, "masked_lm_accuracy": avg_lm_acc})

      # save model if achieved target
      if not achieved and avg_lm_acc >= target:
        wc_end = time.perf_counter()
-        if not os.path.exists(ckpt_dir := save_ckpt_dir): os.mkdir(ckpt_dir)
-        fn = f"{ckpt_dir}/bert-large.safe"
-        safe_save(get_state_dict(model), fn)
-        print(f" *** Model saved to {fn} ***")
+        if getenv("CKPT"):
+          if not os.path.exists(ckpt_dir := save_ckpt_dir): os.mkdir(ckpt_dir)
+          fn = f"{ckpt_dir}/bert-large.safe"
+          safe_save(get_state_dict(model), fn)
+          print(f" *** Model saved to {fn} ***")

        total_seconds = wc_end - wc_start
        hours = int(total_seconds // 3600)
@@ -651,11 +877,12 @@ def train_bert():
        print(f"Reference Convergence point reached after {i * BS} datasamples and {hours}h{minutes}m{seconds:.2f}s.")
        achieved = True
        if MLLOGGER and RUNMLPERF:
+          MLLOGGER.event(key=mllog_constants.EPOCH_STOP, value=i*BS, metadata={"epoch_num": i*BS})
          MLLOGGER.end(key=mllog_constants.RUN_STOP, metadata=dict(status=mllog_constants.SUCCESS))
        # stop once hitting the target
        break

-    if getenv("CKPT", 1) and i % save_ckpt_freq == 0:
+    if getenv("CKPT") and i % save_ckpt_freq == 0:
      if MLLOGGER and RUNMLPERF:
        if previous_step:
          MLLOGGER.end(key=mllog_constants.BLOCK_STOP, value=None, metadata={"first_epoch_num": 1, "epoch_num": 1, "first_step_num": i, "step_num": i, "step_count": i - previous_step})
@@ -0,0 +1,19 @@
+#!/bin/bash
+
+if [ -z $BASEDIR ]; then
+  export BASEDIR="./extra/datasets/"
+fi
+
+cd $BASEDIR
+if [ -d "kits19" ]; then
+  echo "kits19 dataset is already available"
+else
+  echo "Downloading and preparing kits19 dataset at $BASEDIR"
+
+  git clone https://github.com/neheller/kits19
+  cd kits19
+  pip3 install -r requirements.txt
+  python3 -m starter_code.get_imaging
+
+  echo "Done"
+fi
@@ -19,19 +19,8 @@ Install the p2p driver per [README](https://github.com/tinygrad/open-gpu-kernel-
 This is the default on production tinybox green.

 ### tinybox_red
-Disable cwsr
-This is the default on production tinybox red.
-```
-sudo vi /etc/modprobe.d/amdgpu.conf
-cat <<EOF > /etc/modprobe.d/amdgpu.conf
-options amdgpu cwsr_enable=0
-EOF
-sudo update-initramfs -u
-sudo reboot
-
-# validate
-sudo cat /sys/module/amdgpu/parameters/cwsr_enable #= 0
-```
+Disable cwsr + increase mes timeout.
+Install the custom amdgpu driver per [README](https://github.com/nimlgen/amdgpu_ubuntu_22_04/blob/v6.1.3/readme.md)

 # 2. Directions

@@ -4,12 +4,10 @@ export PYTHONPATH="."
 export MODEL="bert"
 export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=4 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=512
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
-
 export BENCHMARK=10 DEBUG=2

 python3 examples/mlperf/model_train.py
@@ -4,12 +4,10 @@ export PYTHONPATH="."
 export MODEL="bert"
 export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=4 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=512
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
+export WANDB=1 PARALLEL=0

-export WANDB=1
-
-python3 examples/mlperf/model_train.py
+RUNMLPERF=1 python3 examples/mlperf/model_train.py
@@ -5,12 +5,10 @@ export MODEL="bert"
 export SUBMISSION_PLATFORM="tinybox_green"
 export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=4 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=512
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
-
 # pip install -e ".[mlperf]"
 export LOGMLPERF=1

@@ -22,4 +20,4 @@ LOGFILE="bert_green_${DATETIME}_${SEED}.log"
 BENCHMARK=10 INITMLPERF=1 python3 examples/mlperf/model_train.py | tee $LOGFILE

 # run
-RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
+PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
@@ -19,19 +19,8 @@ Install the p2p driver per [README](https://github.com/tinygrad/open-gpu-kernel-
 This is the default on production tinybox green.

 ### tinybox_red
-Disable cwsr
-This is the default on production tinybox red.
-```
-sudo vi /etc/modprobe.d/amdgpu.conf
-cat <<EOF > /etc/modprobe.d/amdgpu.conf
-options amdgpu cwsr_enable=0
-EOF
-sudo update-initramfs -u
-sudo reboot
-
-# validate
-sudo cat /sys/module/amdgpu/parameters/cwsr_enable #= 0
-```
+Disable cwsr + increase mes timeout.
+Install the custom amdgpu driver per [README](https://github.com/nimlgen/amdgpu_ubuntu_22_04/blob/v6.1.3/readme.md)

 # 2. Directions

@@ -2,14 +2,12 @@

 export PYTHONPATH="."
 export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=84 EVAL_BS=6
+export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=3
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
-
 export BENCHMARK=10 DEBUG=2

 python3 examples/mlperf/model_train.py
@@ -2,14 +2,12 @@

 export PYTHONPATH="."
 export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=84 EVAL_BS=6
+export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=3
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
+export WANDB=1 PARALLEL=0

-export WANDB=1
-
-python3 examples/mlperf/model_train.py
+RUNMLPERF=1 python3 examples/mlperf/model_train.py
@@ -3,23 +3,21 @@
 export PYTHONPATH="."
 export MODEL="bert"
 export SUBMISSION_PLATFORM="tinybox_red"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=84 EVAL_BS=6
+export DEFAULT_FLOAT="HALF" GPUS=6 BS=66 EVAL_BS=6

-export BEAM=4
+export BEAM=3
+export IGNORE_JIT_FIRST_BEAM=1
 export BASEDIR="/raid/datasets/wiki"

-echo "TODO: DISABLING DROPOUT - UNSET FOR REAL SUBMISSION RUN"
-export DISABLE_DROPOUT=1 # TODO: Unset flag for real submission run.
-
 # pip install -e ".[mlperf]"
 export LOGMLPERF=1

 export SEED=$RANDOM
 DATETIME=$(date "+%m%d%H%M")
-LOGFILE="bert_green_${DATETIME}_${SEED}.log"
+LOGFILE="bert_red_${DATETIME}_${SEED}.log"

 # init
 BENCHMARK=10 INITMLPERF=1 python3 examples/mlperf/model_train.py | tee $LOGFILE

 # run
-RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
+PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
@@ -5,4 +5,4 @@ rocm-smi --setmclk 3
 rocm-smi --setperflevel high

 # power cap to 350W
-echo "350000000" | sudo tee /sys/class/drm/card{1..6}/device/hwmon/hwmon*/power1_cap
+# echo "350000000" | sudo tee /sys/class/drm/card{1..6}/device/hwmon/hwmon*/power1_cap
@@ -1,50 +0,0 @@
-# 1. Problem
-
-This problem uses the ResNet-50 CNN to do image classification.
-
-## Requirements
-
-Install tinygrad and mlperf-logging from master.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-
-### tinybox_green
-Install the p2p driver per [README](https://github.com/tinygrad/open-gpu-kernel-modules/blob/550.54.15-p2p/README.md)
-This is the default on production tinybox green.
-
-### tinybox_red
-Disable cwsr
-This is the default on production tinybox red.
-```
-sudo vi /etc/modprobe.d/amdgpu.conf
-cat <<EOF > /etc/modprobe.d/amdgpu.conf
-options amdgpu cwsr_enable=0
-EOF
-sudo update-initramfs -u
-sudo reboot
-
-# validate
-sudo cat /sys/module/amdgpu/parameters/cwsr_enable #= 0
-```
-
-# 2. Directions
-
-## Steps to download and verify data
-
-```
-IMGNET_TRAIN=1 python3 extra/datasets/imagenet_download.py
-```
-
-## Steps for one time setup
-
-### tinybox_red
-```
-examples/mlperf/training_submission_v4.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/setup.sh
-```
-
-## Steps to run benchmark
-```
-examples/mlperf/training_submission_v4.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/run_and_time.sh
-```
@@ -1,13 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=10 BEAM_PADTO=0
-
-export BENCHMARK=10 DEBUG=2
-
-python3 examples/mlperf/model_train.py
@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=10 BEAM_PADTO=0
-
-export EVAL_START_EPOCH=3 EVAL_FREQ=4
-
-export WANDB=1 PARALLEL=0
-
-python3 examples/mlperf/model_train.py
@@ -1,23 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export SUBMISSION_PLATFORM="tinybox_green"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=10 BEAM_PADTO=0
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="resnet_green_${DATETIME}_${SEED}.log"
-
-# init
-BENCHMARK=10 INITMLPERF=1 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 EVAL_START_EPOCH=3 EVAL_FREQ=4 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
@@ -1,50 +0,0 @@
-# 1. Problem
-
-This problem uses the ResNet-50 CNN to do image classification.
-
-## Requirements
-
-Install tinygrad and mlperf-logging from master.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-
-### tinybox_green
-Install the p2p driver per [README](https://github.com/tinygrad/open-gpu-kernel-modules/blob/550.54.15-p2p/README.md)
-This is the default on production tinybox green.
-
-### tinybox_red
-Disable cwsr
-This is the default on production tinybox red.
-```
-sudo vi /etc/modprobe.d/amdgpu.conf
-cat <<EOF > /etc/modprobe.d/amdgpu.conf
-options amdgpu cwsr_enable=0
-EOF
-sudo update-initramfs -u
-sudo reboot
-
-# validate
-sudo cat /sys/module/amdgpu/parameters/cwsr_enable #= 0
-```
-
-# 2. Directions
-
-## Steps to download and verify data
-
-```
-IMGNET_TRAIN=1 python3 extra/datasets/imagenet_download.py
-```
-
-## Steps for one time setup
-
-### tinybox_red
-```
-examples/mlperf/training_submission_v4.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/setup.sh
-```
-
-## Steps to run benchmark
-```
-examples/mlperf/training_submission_v4.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/run_and_time.sh
-```
@@ -1,13 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=96 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export BENCHMARK=10 DEBUG=2
-
-python3 examples/mlperf/model_train.py
@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=96 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export EVAL_START_EPOCH=3 EVAL_FREQ=4
-
-export WANDB=1 PARALLEL=0
-
-python3 examples/mlperf/model_train.py
@@ -1,23 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="."
-export MODEL="resnet"
-export SUBMISSION_PLATFORM="tinybox_red"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export LAZYCACHE=0 RESET_STEP=0
-
-export TRAIN_BEAM=4 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=2000 BEAM_UPCAST_MAX=96 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="resnet_red_${DATETIME}_${SEED}.log"
-
-# init
-BENCHMARK=10 INITMLPERF=1 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 EVAL_START_EPOCH=3 EVAL_FREQ=4 python3 examples/mlperf/model_train.py | tee -a $LOGFILE
@@ -1,8 +0,0 @@
-#!/bin/bash
-
-rocm-smi --setprofile compute
-rocm-smi --setmclk 3
-rocm-smi --setperflevel high
-
-# power cap to 350W
-echo "350000000" | sudo tee /sys/class/drm/card{1..6}/device/hwmon/hwmon*/power1_cap
@@ -0,0 +1,87 @@
+:::MLLOG {"namespace": "", "time_ms": 1728516968768, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968782, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968782, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968782, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968782, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968917, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728516968917, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728518095273, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728518110874, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125400, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125401, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125401, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125401, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125401, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125401, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125402, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125403, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125403, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728518125403, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728518171154, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728519204577, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728519263743, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728519263744, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38641827217854635, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38641827217854635}}
+:::MLLOG {"namespace": "", "time_ms": 1728520267792, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728520321266, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728520321266, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40444660376272445, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40444660376272445}}
+:::MLLOG {"namespace": "", "time_ms": 1728521322547, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728521376298, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728521376298, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4533385156548231, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4533385156548231}}
+:::MLLOG {"namespace": "", "time_ms": 1728522377080, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728522429361, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728522429362, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5267527467952778, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5267527467952778}}
+:::MLLOG {"namespace": "", "time_ms": 1728523431856, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728523485950, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728523485951, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6413663900499224, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6413663900499224}}
+:::MLLOG {"namespace": "", "time_ms": 1728524487365, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728524539365, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728524539365, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7015928945715869, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.7015928945715869}}
+:::MLLOG {"namespace": "", "time_ms": 1728525540578, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728525593684, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728525593684, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7095211997458373, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7095211997458373}}
+:::MLLOG {"namespace": "", "time_ms": 1728526596673, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728526649935, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728526649936, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.712105579231768, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.712105579231768}}
+:::MLLOG {"namespace": "", "time_ms": 1728527652974, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728527704774, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728527704774, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7137153520152179, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7137153520152179}}
+:::MLLOG {"namespace": "", "time_ms": 1728528706518, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728528759227, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728528759227, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7148766237672532, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7148766237672532}}
+:::MLLOG {"namespace": "", "time_ms": 1728529761683, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728529814822, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728529814823, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7160955339258992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7160955339258992}}
+:::MLLOG {"namespace": "", "time_ms": 1728530822786, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728530876907, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728530876907, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716297444534931, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.716297444534931}}
+:::MLLOG {"namespace": "", "time_ms": 1728531883583, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728531936703, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728531936703, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7177683101775908, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7177683101775908}}
+:::MLLOG {"namespace": "", "time_ms": 1728532940983, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728532993183, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728532993183, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718530326026889, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.718530326026889}}
+:::MLLOG {"namespace": "", "time_ms": 1728534011810, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728534065533, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728534065533, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7185927641985298, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7185927641985298}}
+:::MLLOG {"namespace": "", "time_ms": 1728535067978, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728535122144, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728535122145, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192297569276619, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7192297569276619}}
+:::MLLOG {"namespace": "", "time_ms": 1728536131543, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728536184105, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728536184105, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7197693878473032, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7197693878473032}}
+:::MLLOG {"namespace": "", "time_ms": 1728537200158, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728537253408, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728537253408, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7199979490266993, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7199979490266993}}
+:::MLLOG {"namespace": "", "time_ms": 1728538262628, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728538316185, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728538316186, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7208575420416825, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7208575420416825}}
+:::MLLOG {"namespace": "", "time_ms": 1728538316186, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2849088}}
+:::MLLOG {"namespace": "", "time_ms": 1728538316186, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728538316186, "event_type": "POINT_IN_TIME", "key": "seed", "value": 6505, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,99 @@
+:::MLLOG {"namespace": "", "time_ms": 1728538334148, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334162, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334162, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334162, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334162, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334302, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728538334302, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728539452588, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728539466234, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480889, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480890, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480891, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480892, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728539480892, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728539531181, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728540563757, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728540627488, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728540627488, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38798193001575504, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38798193001575504}}
+:::MLLOG {"namespace": "", "time_ms": 1728541627605, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728541684840, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728541684840, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4043695551053306, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4043695551053306}}
+:::MLLOG {"namespace": "", "time_ms": 1728542682448, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728542739985, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728542739985, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.43848311595381845, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.43848311595381845}}
+:::MLLOG {"namespace": "", "time_ms": 1728543736385, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728543792735, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728543792736, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.519008471802029, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.519008471802029}}
+:::MLLOG {"namespace": "", "time_ms": 1728544788834, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728544846253, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728544846253, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6122590443583112, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6122590443583112}}
+:::MLLOG {"namespace": "", "time_ms": 1728545843895, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728545901605, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728545901606, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.687927868134545, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.687927868134545}}
+:::MLLOG {"namespace": "", "time_ms": 1728546896899, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728546954045, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728546954045, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7037530967627161, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7037530967627161}}
+:::MLLOG {"namespace": "", "time_ms": 1728547951266, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728548007399, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728548007400, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7094319296154922, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7094319296154922}}
+:::MLLOG {"namespace": "", "time_ms": 1728549002681, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728549061015, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728549061015, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7113743569225913, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7113743569225913}}
+:::MLLOG {"namespace": "", "time_ms": 1728550055497, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728550112524, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728550112524, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7128702056715427, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7128702056715427}}
+:::MLLOG {"namespace": "", "time_ms": 1728551106231, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728551163221, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728551163221, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7141216593941458, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7141216593941458}}
+:::MLLOG {"namespace": "", "time_ms": 1728552166923, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728552223146, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728552223147, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7146758918069978, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7146758918069978}}
+:::MLLOG {"namespace": "", "time_ms": 1728553218362, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728553275684, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728553275684, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7159770885197503, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7159770885197503}}
+:::MLLOG {"namespace": "", "time_ms": 1728554271310, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728554328503, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728554328503, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7166727805252052, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7166727805252052}}
+:::MLLOG {"namespace": "", "time_ms": 1728555323857, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728555381026, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728555381026, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7173972926457342, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7173972926457342}}
+:::MLLOG {"namespace": "", "time_ms": 1728556388293, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728556444413, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728556444413, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175414781407389, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7175414781407389}}
+:::MLLOG {"namespace": "", "time_ms": 1728557446881, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728557503528, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728557503528, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7179403612909735, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7179403612909735}}
+:::MLLOG {"namespace": "", "time_ms": 1728558506474, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728558562665, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728558562665, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7183986956585505, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7183986956585505}}
+:::MLLOG {"namespace": "", "time_ms": 1728559573313, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728559631801, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728559631801, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190847733311119, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7190847733311119}}
+:::MLLOG {"namespace": "", "time_ms": 1728560634330, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728560690618, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728560690618, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7193026458280274, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7193026458280274}}
+:::MLLOG {"namespace": "", "time_ms": 1728561686592, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728561743999, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728561743999, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195320777644207, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7195320777644207}}
+:::MLLOG {"namespace": "", "time_ms": 1728562740953, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728562797084, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728562797084, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198491232010441, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.7198491232010441}}
+:::MLLOG {"namespace": "", "time_ms": 1728563794843, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3448896, "step_num": 52256}}
+:::MLLOG {"namespace": "", "time_ms": 1728563851217, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3448896, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3448896, "step_num": 52256, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728563851218, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7201704975009752, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3448896, "masked_lm_accuracy": 0.7201704975009752}}
+:::MLLOG {"namespace": "", "time_ms": 1728563851218, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3448896, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3448896}}
+:::MLLOG {"namespace": "", "time_ms": 1728563851218, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728563851218, "event_type": "POINT_IN_TIME", "key": "seed", "value": 20151, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,100 @@
+:::MLLOG {"namespace": "", "time_ms": 1728563867610, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867623, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867623, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867623, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867623, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867760, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728563867761, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728564959716, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728564973303, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988006, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988006, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988006, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988006, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988006, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988007, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988008, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988008, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988008, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988008, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728564988008, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728565046616, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728566082579, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728566146228, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728566146228, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3875764702688668, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3875764702688668}}
+:::MLLOG {"namespace": "", "time_ms": 1728567148838, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728567205854, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728567205855, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4064989212786429, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4064989212786429}}
+:::MLLOG {"namespace": "", "time_ms": 1728568204176, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728568261886, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728568261886, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4445486674068499, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4445486674068499}}
+:::MLLOG {"namespace": "", "time_ms": 1728569260373, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728569317827, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728569317827, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5039052128362741, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5039052128362741}}
+:::MLLOG {"namespace": "", "time_ms": 1728570316014, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728570373692, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728570373692, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5860576768012982, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5860576768012982}}
+:::MLLOG {"namespace": "", "time_ms": 1728571372081, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728571429048, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728571429048, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6553790274678981, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6553790274678981}}
+:::MLLOG {"namespace": "", "time_ms": 1728572429036, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728572485646, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728572485646, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.695299510013292, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.695299510013292}}
+:::MLLOG {"namespace": "", "time_ms": 1728573485353, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728573541985, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728573541985, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7060198057319994, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7060198057319994}}
+:::MLLOG {"namespace": "", "time_ms": 1728574539420, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728574596892, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728574596892, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7096440969074137, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7096440969074137}}
+:::MLLOG {"namespace": "", "time_ms": 1728575594573, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728575651826, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728575651827, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7117632173033053, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7117632173033053}}
+:::MLLOG {"namespace": "", "time_ms": 1728576649028, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728576706272, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728576706273, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7133303043652096, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7133303043652096}}
+:::MLLOG {"namespace": "", "time_ms": 1728577704079, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728577761356, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728577761356, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7137490666835123, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7137490666835123}}
+:::MLLOG {"namespace": "", "time_ms": 1728578758647, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728578815837, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728578815837, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7155028079610137, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7155028079610137}}
+:::MLLOG {"namespace": "", "time_ms": 1728579812964, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728579870775, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728579870776, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7161324361018528, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7161324361018528}}
+:::MLLOG {"namespace": "", "time_ms": 1728580866942, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728580924299, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728580924299, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7167852242763842, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7167852242763842}}
+:::MLLOG {"namespace": "", "time_ms": 1728581922472, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728581979021, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728581979022, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7176919813705335, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7176919813705335}}
+:::MLLOG {"namespace": "", "time_ms": 1728582983990, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728583040463, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728583040464, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7181049583197069, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7181049583197069}}
+:::MLLOG {"namespace": "", "time_ms": 1728584044016, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728584101869, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728584101869, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7182932556855443, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7182932556855443}}
+:::MLLOG {"namespace": "", "time_ms": 1728585118787, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728585176571, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728585176571, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189288123968338, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7189288123968338}}
+:::MLLOG {"namespace": "", "time_ms": 1728586176081, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728586233843, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728586233843, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190464254475384, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7190464254475384}}
+:::MLLOG {"namespace": "", "time_ms": 1728587233422, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728587290110, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728587290111, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189439674587971, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7189439674587971}}
+:::MLLOG {"namespace": "", "time_ms": 1728588298920, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728588356857, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728588356858, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195348424974429, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.7195348424974429}}
+:::MLLOG {"namespace": "", "time_ms": 1728589362222, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3448896, "step_num": 52256}}
+:::MLLOG {"namespace": "", "time_ms": 1728589419360, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3448896, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3448896, "step_num": 52256, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728589419360, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198157412389402, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3448896, "masked_lm_accuracy": 0.7198157412389402}}
+:::MLLOG {"namespace": "", "time_ms": 1728590419216, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3598848, "step_num": 54528}}
+:::MLLOG {"namespace": "", "time_ms": 1728590477169, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3598848, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3598848, "step_num": 54528, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728590477170, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7196774663650568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3598848, "masked_lm_accuracy": 0.7196774663650568}}
+:::MLLOG {"namespace": "", "time_ms": 1728590477170, "event_type": "POINT_IN_TIME", "key": "seed", "value": 15936, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,78 @@
+:::MLLOG {"namespace": "", "time_ms": 1728590796935, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796948, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796948, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796949, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796949, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728590797097, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728590797098, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728591915361, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728591929057, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943823, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943823, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943826, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728591989241, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728593032285, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728593096400, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728593096400, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38762320266130373, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38762320266130373}}
+:::MLLOG {"namespace": "", "time_ms": 1728594107968, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728594165957, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728594165958, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.41101139415576204, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.41101139415576204}}
+:::MLLOG {"namespace": "", "time_ms": 1728595175795, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728595232491, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728595232492, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4496020218106037, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4496020218106037}}
+:::MLLOG {"namespace": "", "time_ms": 1728596242699, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728596299316, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728596299317, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5736796348911599, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5736796348911599}}
+:::MLLOG {"namespace": "", "time_ms": 1728597307215, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728597365230, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728597365230, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6932630223218166, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6932630223218166}}
+:::MLLOG {"namespace": "", "time_ms": 1728598375878, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728598433374, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728598433374, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7074951962503617, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.7074951962503617}}
+:::MLLOG {"namespace": "", "time_ms": 1728599442119, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728599498818, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728599498818, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7122511012366809, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7122511012366809}}
+:::MLLOG {"namespace": "", "time_ms": 1728600507693, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728600566544, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728600566545, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7142625644525941, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7142625644525941}}
+:::MLLOG {"namespace": "", "time_ms": 1728601572961, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728601630426, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728601630427, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7150282886618973, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7150282886618973}}
+:::MLLOG {"namespace": "", "time_ms": 1728602641519, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728602697877, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728602697878, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7163425168378953, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7163425168378953}}
+:::MLLOG {"namespace": "", "time_ms": 1728603708872, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728603765541, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728603765541, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7169951724305293, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7169951724305293}}
+:::MLLOG {"namespace": "", "time_ms": 1728604783436, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728604840983, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728604840983, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7174601919220533, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7174601919220533}}
+:::MLLOG {"namespace": "", "time_ms": 1728605853878, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728605911397, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728605911397, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718169204153268, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.718169204153268}}
+:::MLLOG {"namespace": "", "time_ms": 1728606921863, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728606978402, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728606978403, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192717549253096, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7192717549253096}}
+:::MLLOG {"namespace": "", "time_ms": 1728607995812, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728608053682, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728608053682, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7194449275499629, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7194449275499629}}
+:::MLLOG {"namespace": "", "time_ms": 1728609084329, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142724, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200488402375792, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7200488402375792}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2399232}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "seed", "value": 20762, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,96 @@
+:::MLLOG {"namespace": "", "time_ms": 1728609159110, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159123, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159123, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159124, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159124, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159270, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728609159271, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728610267217, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728610281092, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295781, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295781, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295781, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295782, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728610295783, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728610347873, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728611397047, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728611461353, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728611461353, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38775324024121494, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38775324024121494}}
+:::MLLOG {"namespace": "", "time_ms": 1728612481307, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728612538833, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728612538833, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4005705962727437, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4005705962727437}}
+:::MLLOG {"namespace": "", "time_ms": 1728613556463, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728613613480, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728613613480, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4705434759434069, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4705434759434069}}
+:::MLLOG {"namespace": "", "time_ms": 1728614631311, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728614687851, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728614687851, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5429764502622013, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5429764502622013}}
+:::MLLOG {"namespace": "", "time_ms": 1728615706690, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728615763381, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728615763381, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6442740782693109, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6442740782693109}}
+:::MLLOG {"namespace": "", "time_ms": 1728616781291, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728616838087, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728616838087, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6980952795351345, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6980952795351345}}
+:::MLLOG {"namespace": "", "time_ms": 1728617855821, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728617912283, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728617912283, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.707689059052413, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.707689059052413}}
+:::MLLOG {"namespace": "", "time_ms": 1728618929847, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728618986661, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728618986661, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7109381170707616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7109381170707616}}
+:::MLLOG {"namespace": "", "time_ms": 1728620004644, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728620061231, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728620061231, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.712553325461712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.712553325461712}}
+:::MLLOG {"namespace": "", "time_ms": 1728621080635, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728621138171, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728621138171, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.714370012497859, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.714370012497859}}
+:::MLLOG {"namespace": "", "time_ms": 1728622165561, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728622221929, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728622221929, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7153810627029982, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7153810627029982}}
+:::MLLOG {"namespace": "", "time_ms": 1728623242362, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728623299587, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728623299588, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7157506428582981, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7157506428582981}}
+:::MLLOG {"namespace": "", "time_ms": 1728624319084, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728624376686, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728624376687, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716899270201845, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.716899270201845}}
+:::MLLOG {"namespace": "", "time_ms": 1728625402109, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728625460126, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728625460126, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7170993249861152, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7170993249861152}}
+:::MLLOG {"namespace": "", "time_ms": 1728626480567, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728626539447, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728626539448, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7173656942820077, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7173656942820077}}
+:::MLLOG {"namespace": "", "time_ms": 1728627566869, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728627623861, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728627623862, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7182765536917565, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7182765536917565}}
+:::MLLOG {"namespace": "", "time_ms": 1728628643742, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728628701681, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728628701681, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7185246580435118, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7185246580435118}}
+:::MLLOG {"namespace": "", "time_ms": 1728629728019, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728629785903, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728629785903, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7187082183954597, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7187082183954597}}
+:::MLLOG {"namespace": "", "time_ms": 1728630806207, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728630863277, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728630863277, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192723798623111, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7192723798623111}}
+:::MLLOG {"namespace": "", "time_ms": 1728631896497, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728631954615, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728631954616, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195008011299047, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7195008011299047}}
+:::MLLOG {"namespace": "", "time_ms": 1728632984222, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728633041346, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728633041346, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195586102339202, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7195586102339202}}
+:::MLLOG {"namespace": "", "time_ms": 1728634068369, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728634127804, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728634127805, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7202416826810534, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.7202416826810534}}
+:::MLLOG {"namespace": "", "time_ms": 1728634127805, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3298944}}
+:::MLLOG {"namespace": "", "time_ms": 1728634127805, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728634127805, "event_type": "POINT_IN_TIME", "key": "seed", "value": 219, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,93 @@
+:::MLLOG {"namespace": "", "time_ms": 1728524965239, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965253, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965253, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965253, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965253, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965392, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728524965393, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728526062933, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728526077625, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092468, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092468, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092468, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092468, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092469, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728526092470, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728526138852, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728527173955, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728527237057, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728527237057, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3877335678748049, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3877335678748049}}
+:::MLLOG {"namespace": "", "time_ms": 1728528242671, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728528300179, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728528300179, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40263391588716785, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40263391588716785}}
+:::MLLOG {"namespace": "", "time_ms": 1728529303573, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728529359959, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728529359959, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.42964756483436706, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.42964756483436706}}
+:::MLLOG {"namespace": "", "time_ms": 1728530364689, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728530422126, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728530422126, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4922114010334873, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.4922114010334873}}
+:::MLLOG {"namespace": "", "time_ms": 1728531424135, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728531481285, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728531481286, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5822947872707639, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5822947872707639}}
+:::MLLOG {"namespace": "", "time_ms": 1728532482693, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728532539641, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728532539641, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6813526992749224, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6813526992749224}}
+:::MLLOG {"namespace": "", "time_ms": 1728533542012, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728533599055, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728533599056, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7051245396052854, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7051245396052854}}
+:::MLLOG {"namespace": "", "time_ms": 1728534601465, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728534657337, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728534657337, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7107940144930761, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7107940144930761}}
+:::MLLOG {"namespace": "", "time_ms": 1728535659894, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728535716827, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728535716827, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7130225877169728, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7130225877169728}}
+:::MLLOG {"namespace": "", "time_ms": 1728536719704, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728536776465, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728536776465, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7149506402573474, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7149506402573474}}
+:::MLLOG {"namespace": "", "time_ms": 1728537786290, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728537844483, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728537844483, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7156235970418183, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7156235970418183}}
+:::MLLOG {"namespace": "", "time_ms": 1728538854316, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728538910323, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728538910323, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716550941694691, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.716550941694691}}
+:::MLLOG {"namespace": "", "time_ms": 1728539924500, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728539982702, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728539982702, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7173966967065533, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7173966967065533}}
+:::MLLOG {"namespace": "", "time_ms": 1728540985244, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728541042328, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728541042329, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7179239220295971, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7179239220295971}}
+:::MLLOG {"namespace": "", "time_ms": 1728542045079, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728542102249, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728542102249, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7180614758195746, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7180614758195746}}
+:::MLLOG {"namespace": "", "time_ms": 1728543106186, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728543163150, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728543163150, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7188610007967431, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7188610007967431}}
+:::MLLOG {"namespace": "", "time_ms": 1728544166236, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728544222285, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728544222285, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190189941528677, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7190189941528677}}
+:::MLLOG {"namespace": "", "time_ms": 1728545224291, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728545280273, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728545280273, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189939534442469, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7189939534442469}}
+:::MLLOG {"namespace": "", "time_ms": 1728546290997, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728546348255, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728546348255, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7197046605641069, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7197046605641069}}
+:::MLLOG {"namespace": "", "time_ms": 1728547370392, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728547427703, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728547427704, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198148297920296, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7198148297920296}}
+:::MLLOG {"namespace": "", "time_ms": 1728548433236, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728548489162, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728548489163, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7201927967892483, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7201927967892483}}
+:::MLLOG {"namespace": "", "time_ms": 1728548489163, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3148992}}
+:::MLLOG {"namespace": "", "time_ms": 1728548489163, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728548489163, "event_type": "POINT_IN_TIME", "key": "seed", "value": 28210, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,87 @@
+:::MLLOG {"namespace": "", "time_ms": 1728548506444, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506457, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506458, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506458, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506458, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506735, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728548506735, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728549621641, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728549635270, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652872, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652872, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652872, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652872, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652872, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652873, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652874, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652874, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728549652874, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728549705049, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728550749936, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728550814006, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728550814007, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3879261584913366, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3879261584913366}}
+:::MLLOG {"namespace": "", "time_ms": 1728551828200, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728551884657, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728551884657, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40485212616242544, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40485212616242544}}
+:::MLLOG {"namespace": "", "time_ms": 1728552894731, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728552950982, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728552950982, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.44991275376735795, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.44991275376735795}}
+:::MLLOG {"namespace": "", "time_ms": 1728553961756, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728554019767, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728554019768, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5169473737210089, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5169473737210089}}
+:::MLLOG {"namespace": "", "time_ms": 1728555028196, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728555084535, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728555084536, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6089998500284207, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6089998500284207}}
+:::MLLOG {"namespace": "", "time_ms": 1728556093244, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728556150861, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728556150861, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6900064689568152, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6900064689568152}}
+:::MLLOG {"namespace": "", "time_ms": 1728557159270, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728557216649, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728557216649, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7058100798706416, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7058100798706416}}
+:::MLLOG {"namespace": "", "time_ms": 1728558223895, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728558281962, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728558281962, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7103145412148726, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7103145412148726}}
+:::MLLOG {"namespace": "", "time_ms": 1728559289692, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728559346801, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728559346801, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7119302401016341, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7119302401016341}}
+:::MLLOG {"namespace": "", "time_ms": 1728560353303, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728560410326, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728560410327, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7139668108176956, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7139668108176956}}
+:::MLLOG {"namespace": "", "time_ms": 1728561418512, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728561474810, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728561474810, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.715211286184383, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.715211286184383}}
+:::MLLOG {"namespace": "", "time_ms": 1728562483778, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728562541169, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728562541169, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7157247102968551, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7157247102968551}}
+:::MLLOG {"namespace": "", "time_ms": 1728563551560, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728563609899, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728563609899, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.717246541128805, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.717246541128805}}
+:::MLLOG {"namespace": "", "time_ms": 1728564615944, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728564673046, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728564673046, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7174253223824801, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7174253223824801}}
+:::MLLOG {"namespace": "", "time_ms": 1728565680391, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728565737965, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728565737965, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175353851658753, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7175353851658753}}
+:::MLLOG {"namespace": "", "time_ms": 1728566749670, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728566807129, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728566807130, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7185968662280842, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7185968662280842}}
+:::MLLOG {"namespace": "", "time_ms": 1728567827685, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728567883857, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728567883857, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189708248755141, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7189708248755141}}
+:::MLLOG {"namespace": "", "time_ms": 1728568903914, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728568961416, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728568961416, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7191977317131559, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7191977317131559}}
+:::MLLOG {"namespace": "", "time_ms": 1728569970441, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728570027956, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728570027956, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7201856829552287, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7201856829552287}}
+:::MLLOG {"namespace": "", "time_ms": 1728570027957, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2849088}}
+:::MLLOG {"namespace": "", "time_ms": 1728570027957, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728570027957, "event_type": "POINT_IN_TIME", "key": "seed", "value": 10448, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,96 @@
+:::MLLOG {"namespace": "", "time_ms": 1728570044042, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044056, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044056, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044056, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044056, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044375, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728570044375, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728571155314, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728571169188, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185077, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185077, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185077, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185077, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185077, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185078, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185079, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185079, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185079, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185079, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728571185079, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728571232077, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728572260932, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728572324272, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728572324272, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3880780045687258, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3880780045687258}}
+:::MLLOG {"namespace": "", "time_ms": 1728573321862, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728573379098, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728573379099, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4047361867257629, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4047361867257629}}
+:::MLLOG {"namespace": "", "time_ms": 1728574372317, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728574430334, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728574430334, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4497806984254108, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4497806984254108}}
+:::MLLOG {"namespace": "", "time_ms": 1728575423648, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728575480052, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728575480052, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5134892886673252, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5134892886673252}}
+:::MLLOG {"namespace": "", "time_ms": 1728576473686, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728576529490, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728576529490, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6072650871808899, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6072650871808899}}
+:::MLLOG {"namespace": "", "time_ms": 1728577521455, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728577578395, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728577578395, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6880894238246581, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6880894238246581}}
+:::MLLOG {"namespace": "", "time_ms": 1728578569253, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728578627203, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728578627203, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7046831176176569, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7046831176176569}}
+:::MLLOG {"namespace": "", "time_ms": 1728579618293, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728579675086, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728579675087, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7089016024552924, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7089016024552924}}
+:::MLLOG {"namespace": "", "time_ms": 1728580666335, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728580724803, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728580724804, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7109720907171257, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7109720907171257}}
+:::MLLOG {"namespace": "", "time_ms": 1728581716985, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728581773028, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728581773028, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7131220351884519, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7131220351884519}}
+:::MLLOG {"namespace": "", "time_ms": 1728582764878, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728582822486, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728582822487, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7144350144916047, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7144350144916047}}
+:::MLLOG {"namespace": "", "time_ms": 1728583814857, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728583872551, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728583872551, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7149028261407235, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7149028261407235}}
+:::MLLOG {"namespace": "", "time_ms": 1728584862980, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728584919984, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728584919984, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7159077662798052, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7159077662798052}}
+:::MLLOG {"namespace": "", "time_ms": 1728585920409, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728585976497, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728585976497, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7165625739683987, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7165625739683987}}
+:::MLLOG {"namespace": "", "time_ms": 1728586977974, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728587034924, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728587034924, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716636929886743, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.716636929886743}}
+:::MLLOG {"namespace": "", "time_ms": 1728588046632, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728588103542, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728588103543, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7177939713180029, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7177939713180029}}
+:::MLLOG {"namespace": "", "time_ms": 1728589095819, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728589153039, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728589153040, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718201180382553, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.718201180382553}}
+:::MLLOG {"namespace": "", "time_ms": 1728590145828, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728590203087, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728590203087, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7187695578679255, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7187695578679255}}
+:::MLLOG {"namespace": "", "time_ms": 1728591195213, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728591252415, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728591252415, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192902017941215, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7192902017941215}}
+:::MLLOG {"namespace": "", "time_ms": 1728592251653, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728592309614, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728592309614, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7194759279721928, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7194759279721928}}
+:::MLLOG {"namespace": "", "time_ms": 1728593301651, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728593358684, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728593358684, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7194070329763393, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7194070329763393}}
+:::MLLOG {"namespace": "", "time_ms": 1728594357831, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728594414864, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728594414864, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200078809769052, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.7200078809769052}}
+:::MLLOG {"namespace": "", "time_ms": 1728594414864, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3298944}}
+:::MLLOG {"namespace": "", "time_ms": 1728594414864, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728594414864, "event_type": "POINT_IN_TIME", "key": "seed", "value": 10752, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,78 @@
+:::MLLOG {"namespace": "", "time_ms": 1728590796935, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796948, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796948, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796949, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728590796949, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728590797097, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728590797098, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728591915361, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728591929057, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943823, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943823, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943824, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943825, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728591943826, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728591989241, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728593032285, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728593096400, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728593096400, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38762320266130373, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38762320266130373}}
+:::MLLOG {"namespace": "", "time_ms": 1728594107968, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728594165957, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728594165958, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.41101139415576204, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.41101139415576204}}
+:::MLLOG {"namespace": "", "time_ms": 1728595175795, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728595232491, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728595232492, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4496020218106037, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4496020218106037}}
+:::MLLOG {"namespace": "", "time_ms": 1728596242699, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728596299316, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728596299317, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5736796348911599, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5736796348911599}}
+:::MLLOG {"namespace": "", "time_ms": 1728597307215, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728597365230, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728597365230, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6932630223218166, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6932630223218166}}
+:::MLLOG {"namespace": "", "time_ms": 1728598375878, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728598433374, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728598433374, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7074951962503617, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.7074951962503617}}
+:::MLLOG {"namespace": "", "time_ms": 1728599442119, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728599498818, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728599498818, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7122511012366809, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7122511012366809}}
+:::MLLOG {"namespace": "", "time_ms": 1728600507693, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728600566544, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728600566545, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7142625644525941, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7142625644525941}}
+:::MLLOG {"namespace": "", "time_ms": 1728601572961, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728601630426, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728601630427, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7150282886618973, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7150282886618973}}
+:::MLLOG {"namespace": "", "time_ms": 1728602641519, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728602697877, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728602697878, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7163425168378953, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7163425168378953}}
+:::MLLOG {"namespace": "", "time_ms": 1728603708872, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728603765541, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728603765541, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7169951724305293, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7169951724305293}}
+:::MLLOG {"namespace": "", "time_ms": 1728604783436, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728604840983, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728604840983, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7174601919220533, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7174601919220533}}
+:::MLLOG {"namespace": "", "time_ms": 1728605853878, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728605911397, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728605911397, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718169204153268, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.718169204153268}}
+:::MLLOG {"namespace": "", "time_ms": 1728606921863, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728606978402, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728606978403, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192717549253096, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7192717549253096}}
+:::MLLOG {"namespace": "", "time_ms": 1728607995812, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728608053682, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728608053682, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7194449275499629, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7194449275499629}}
+:::MLLOG {"namespace": "", "time_ms": 1728609084329, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142724, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200488402375792, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7200488402375792}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2399232}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728609142725, "event_type": "POINT_IN_TIME", "key": "seed", "value": 22978, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,90 @@
+:::MLLOG {"namespace": "", "time_ms": 1728618801034, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801047, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_green", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801048, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801048, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801048, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801209, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728618801209, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728619929327, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728619943087, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957912, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957913, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957914, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728619957915, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728620007681, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728621045859, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728621108773, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728621108773, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3865617481387155, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3865617481387155}}
+:::MLLOG {"namespace": "", "time_ms": 1728622114605, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728622171150, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728622171151, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40262117073717557, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40262117073717557}}
+:::MLLOG {"namespace": "", "time_ms": 1728623174052, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728623230873, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728623230873, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4523024549724531, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4523024549724531}}
+:::MLLOG {"namespace": "", "time_ms": 1728624233281, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728624289615, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728624289616, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5221301693185476, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5221301693185476}}
+:::MLLOG {"namespace": "", "time_ms": 1728625291673, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728625349229, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728625349230, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6420224388345102, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6420224388345102}}
+:::MLLOG {"namespace": "", "time_ms": 1728626351043, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728626408667, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728626408667, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.699528387405233, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.699528387405233}}
+:::MLLOG {"namespace": "", "time_ms": 1728627410463, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728627466811, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728627466811, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7079550924241078, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7079550924241078}}
+:::MLLOG {"namespace": "", "time_ms": 1728628468561, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728628525964, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728628525965, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7118257526945195, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7118257526945195}}
+:::MLLOG {"namespace": "", "time_ms": 1728629528193, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728629584413, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728629584413, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7130748185842568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7130748185842568}}
+:::MLLOG {"namespace": "", "time_ms": 1728630587761, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728630645519, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728630645520, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.714902451421661, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.714902451421661}}
+:::MLLOG {"namespace": "", "time_ms": 1728631646942, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728631704444, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728631704445, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7154911232337883, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7154911232337883}}
+:::MLLOG {"namespace": "", "time_ms": 1728632715855, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728632772059, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728632772059, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7161726189646523, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7161726189646523}}
+:::MLLOG {"namespace": "", "time_ms": 1728633776195, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728633832753, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728633832753, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7168811186745844, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7168811186745844}}
+:::MLLOG {"namespace": "", "time_ms": 1728634841112, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728634897563, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728634897563, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7172674153762159, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7172674153762159}}
+:::MLLOG {"namespace": "", "time_ms": 1728635899496, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728635956963, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728635956963, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.717885359040119, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.717885359040119}}
+:::MLLOG {"namespace": "", "time_ms": 1728636964772, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728637021152, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728637021152, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718571519987556, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.718571519987556}}
+:::MLLOG {"namespace": "", "time_ms": 1728638039702, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728638096247, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728638096247, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190522497187993, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7190522497187993}}
+:::MLLOG {"namespace": "", "time_ms": 1728639099786, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728639156384, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728639156384, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190969545610951, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7190969545610951}}
+:::MLLOG {"namespace": "", "time_ms": 1728640172070, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728640229709, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728640229709, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7197579530877272, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7197579530877272}}
+:::MLLOG {"namespace": "", "time_ms": 1728641233638, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728641290427, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728641290427, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200084952587272, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7200084952587272}}
+:::MLLOG {"namespace": "", "time_ms": 1728641290427, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2999040}}
+:::MLLOG {"namespace": "", "time_ms": 1728641290428, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728641290428, "event_type": "POINT_IN_TIME", "key": "seed", "value": 9634, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,93 @@
+:::MLLOG {"namespace": "", "time_ms": 1728516945293, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945306, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945306, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945307, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945307, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945490, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728516945491, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728518502722, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728518514008, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528257, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528257, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528258, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728518528259, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728518578008, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728519789747, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728519848812, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728519848813, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38834677969150794, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38834677969150794}}
+:::MLLOG {"namespace": "", "time_ms": 1728521004546, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728521057409, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728521057409, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40273311821800833, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40273311821800833}}
+:::MLLOG {"namespace": "", "time_ms": 1728522209891, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728522262867, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728522262867, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.44312383646584586, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.44312383646584586}}
+:::MLLOG {"namespace": "", "time_ms": 1728523414970, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728523467811, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728523467812, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5062930049407103, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5062930049407103}}
+:::MLLOG {"namespace": "", "time_ms": 1728524621021, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728524672880, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728524672880, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5888718058230566, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5888718058230566}}
+:::MLLOG {"namespace": "", "time_ms": 1728525825012, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728525878025, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728525878025, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6833749091689574, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6833749091689574}}
+:::MLLOG {"namespace": "", "time_ms": 1728527029872, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728527083655, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728527083655, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7040974767750154, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7040974767750154}}
+:::MLLOG {"namespace": "", "time_ms": 1728528235285, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728528288182, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728528288183, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7093540651038799, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7093540651038799}}
+:::MLLOG {"namespace": "", "time_ms": 1728529439469, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728529492117, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728529492117, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7120359153824791, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7120359153824791}}
+:::MLLOG {"namespace": "", "time_ms": 1728530643295, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728530695998, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728530695999, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7137061499448043, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7137061499448043}}
+:::MLLOG {"namespace": "", "time_ms": 1728531854130, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728531905972, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728531905973, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7148870925168185, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7148870925168185}}
+:::MLLOG {"namespace": "", "time_ms": 1728533058059, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728533110923, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728533110923, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7155579556657943, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7155579556657943}}
+:::MLLOG {"namespace": "", "time_ms": 1728534262752, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728534314582, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728534314582, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716432806969261, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.716432806969261}}
+:::MLLOG {"namespace": "", "time_ms": 1728535466733, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728535518647, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728535518647, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7171081434700685, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7171081434700685}}
+:::MLLOG {"namespace": "", "time_ms": 1728536674344, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728536727210, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728536727211, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175339546901564, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7175339546901564}}
+:::MLLOG {"namespace": "", "time_ms": 1728537884762, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728537937482, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728537937482, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718394822250531, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.718394822250531}}
+:::MLLOG {"namespace": "", "time_ms": 1728539088991, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728539141647, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728539141648, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7188054349298025, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7188054349298025}}
+:::MLLOG {"namespace": "", "time_ms": 1728540309614, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728540362635, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728540362635, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7188799817474859, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7188799817474859}}
+:::MLLOG {"namespace": "", "time_ms": 1728541513809, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728541567687, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728541567687, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195802817962523, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7195802817962523}}
+:::MLLOG {"namespace": "", "time_ms": 1728542725822, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728542778626, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728542778627, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.719959929272595, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.719959929272595}}
+:::MLLOG {"namespace": "", "time_ms": 1728543937336, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728543990349, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728543990349, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200285581392518, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7200285581392518}}
+:::MLLOG {"namespace": "", "time_ms": 1728543990349, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3148992}}
+:::MLLOG {"namespace": "", "time_ms": 1728543990349, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728543990349, "event_type": "POINT_IN_TIME", "key": "seed", "value": 21254, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,99 @@
+:::MLLOG {"namespace": "", "time_ms": 1728544003831, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728544003844, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728544003844, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728544003844, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728544003845, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728544004114, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728544004115, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728545567231, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728545577478, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591839, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591839, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591839, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591839, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591840, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591841, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591841, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591841, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591841, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728545591841, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728545644357, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728546841240, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728546898744, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728546898744, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38783227444481694, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38783227444481694}}
+:::MLLOG {"namespace": "", "time_ms": 1728548039569, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728548092787, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728548092788, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4103626927252508, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4103626927252508}}
+:::MLLOG {"namespace": "", "time_ms": 1728549229970, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728549283151, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728549283152, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.45124479976326815, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.45124479976326815}}
+:::MLLOG {"namespace": "", "time_ms": 1728550420161, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728550473622, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728550473623, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5112407873616508, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5112407873616508}}
+:::MLLOG {"namespace": "", "time_ms": 1728551611312, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728551664679, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728551664679, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6004241028086611, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6004241028086611}}
+:::MLLOG {"namespace": "", "time_ms": 1728552801410, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728552853292, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728552853292, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6789556939443143, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6789556939443143}}
+:::MLLOG {"namespace": "", "time_ms": 1728553989456, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728554042458, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728554042458, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7036405706305524, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7036405706305524}}
+:::MLLOG {"namespace": "", "time_ms": 1728555178691, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728555231909, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728555231909, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7080815071178612, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7080815071178612}}
+:::MLLOG {"namespace": "", "time_ms": 1728556368078, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728556420173, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728556420173, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7108210982072117, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7108210982072117}}
+:::MLLOG {"namespace": "", "time_ms": 1728557556063, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728557608016, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728557608016, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7125928249628013, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7125928249628013}}
+:::MLLOG {"namespace": "", "time_ms": 1728558744467, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728558796348, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728558796348, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7135176017174266, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7135176017174266}}
+:::MLLOG {"namespace": "", "time_ms": 1728559932625, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728559984418, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728559984418, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7145863458576786, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7145863458576786}}
+:::MLLOG {"namespace": "", "time_ms": 1728561120215, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728561173235, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728561173236, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7155747616727265, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7155747616727265}}
+:::MLLOG {"namespace": "", "time_ms": 1728562309022, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728562360931, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728562360931, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7163686522768155, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7163686522768155}}
+:::MLLOG {"namespace": "", "time_ms": 1728563501949, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728563554781, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728563554781, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716544431773359, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.716544431773359}}
+:::MLLOG {"namespace": "", "time_ms": 1728564690440, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728564743201, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728564743201, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175336914214104, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7175336914214104}}
+:::MLLOG {"namespace": "", "time_ms": 1728565883246, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728565936169, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728565936169, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7186442701202992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7186442701202992}}
+:::MLLOG {"namespace": "", "time_ms": 1728567078310, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728567130373, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728567130374, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7184390328784295, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7184390328784295}}
+:::MLLOG {"namespace": "", "time_ms": 1728568270051, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728568322854, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728568322854, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7190133258834455, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7190133258834455}}
+:::MLLOG {"namespace": "", "time_ms": 1728569476727, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728569528906, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728569528906, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7193232258423117, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7193232258423117}}
+:::MLLOG {"namespace": "", "time_ms": 1728570672300, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728570725375, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728570725375, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7193006322136452, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7193006322136452}}
+:::MLLOG {"namespace": "", "time_ms": 1728571862039, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728571914059, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728571914060, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.71975367935961, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.71975367935961}}
+:::MLLOG {"namespace": "", "time_ms": 1728573050479, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3448896, "step_num": 52256}}
+:::MLLOG {"namespace": "", "time_ms": 1728573103486, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3448896, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3448896, "step_num": 52256, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728573103487, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200944615325745, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3448896, "masked_lm_accuracy": 0.7200944615325745}}
+:::MLLOG {"namespace": "", "time_ms": 1728573103487, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3448896, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3448896}}
+:::MLLOG {"namespace": "", "time_ms": 1728573103487, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728573103487, "event_type": "POINT_IN_TIME", "key": "seed", "value": 31023, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,81 @@
+:::MLLOG {"namespace": "", "time_ms": 1728573115586, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115599, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115600, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115600, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115600, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115806, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728573115807, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728574701526, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728574711971, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726511, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726512, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726513, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728574726514, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728574767832, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728575974065, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728576033823, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728576033823, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3871400942148816, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3871400942148816}}
+:::MLLOG {"namespace": "", "time_ms": 1728577185498, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728577237616, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728577237616, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.39147048244665106, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.39147048244665106}}
+:::MLLOG {"namespace": "", "time_ms": 1728578385099, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728578438339, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728578438339, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.45783117753199354, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.45783117753199354}}
+:::MLLOG {"namespace": "", "time_ms": 1728579586799, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728579639018, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728579639018, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5536425763000323, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5536425763000323}}
+:::MLLOG {"namespace": "", "time_ms": 1728580787082, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728580840252, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728580840252, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6808095073800067, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6808095073800067}}
+:::MLLOG {"namespace": "", "time_ms": 1728581987844, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728582039916, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728582039917, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.70571623852529, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.70571623852529}}
+:::MLLOG {"namespace": "", "time_ms": 1728583186452, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728583239391, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728583239392, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7116154495584228, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7116154495584228}}
+:::MLLOG {"namespace": "", "time_ms": 1728584385675, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728584438607, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728584438607, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7138893727993064, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7138893727993064}}
+:::MLLOG {"namespace": "", "time_ms": 1728585586366, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728585638448, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728585638449, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7147087185198344, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7147087185198344}}
+:::MLLOG {"namespace": "", "time_ms": 1728586785128, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728586837097, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728586837097, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7161258058628066, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7161258058628066}}
+:::MLLOG {"namespace": "", "time_ms": 1728587983727, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728588035691, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728588035692, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175371565358254, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7175371565358254}}
+:::MLLOG {"namespace": "", "time_ms": 1728589188060, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728589241204, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728589241204, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.71796901310427, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.71796901310427}}
+:::MLLOG {"namespace": "", "time_ms": 1728590398050, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728590450110, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728590450111, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7184663905522461, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7184663905522461}}
+:::MLLOG {"namespace": "", "time_ms": 1728591601749, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728591654909, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728591654909, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7191987109527519, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7191987109527519}}
+:::MLLOG {"namespace": "", "time_ms": 1728592801533, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728592853668, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728592853668, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7193377415935079, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7193377415935079}}
+:::MLLOG {"namespace": "", "time_ms": 1728594003086, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728594056207, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728594056207, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198944457577029, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7198944457577029}}
+:::MLLOG {"namespace": "", "time_ms": 1728595208245, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728595261648, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728595261648, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7206625282657168, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7206625282657168}}
+:::MLLOG {"namespace": "", "time_ms": 1728595261648, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2549184}}
+:::MLLOG {"namespace": "", "time_ms": 1728595261648, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728595261648, "event_type": "POINT_IN_TIME", "key": "seed", "value": 6117, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,90 @@
+:::MLLOG {"namespace": "", "time_ms": 1728595273658, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273672, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273672, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273672, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273672, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273864, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728595273865, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728596851159, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728596861539, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876123, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876124, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876125, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876126, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876126, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728596876126, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728596921065, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728598131675, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728598191093, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728598191093, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38629304813852405, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38629304813852405}}
+:::MLLOG {"namespace": "", "time_ms": 1728599345629, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728599399811, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728599399812, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4057833703070539, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4057833703070539}}
+:::MLLOG {"namespace": "", "time_ms": 1728600550664, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728600603826, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728600603826, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4394682460404758, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4394682460404758}}
+:::MLLOG {"namespace": "", "time_ms": 1728601755405, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728601807545, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728601807545, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5091635388711576, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5091635388711576}}
+:::MLLOG {"namespace": "", "time_ms": 1728602959507, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728603011549, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728603011550, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6130874210764136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6130874210764136}}
+:::MLLOG {"namespace": "", "time_ms": 1728604163051, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728604215106, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728604215107, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6886828500994252, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6886828500994252}}
+:::MLLOG {"namespace": "", "time_ms": 1728605365869, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728605418812, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728605418812, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7043221683007339, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7043221683007339}}
+:::MLLOG {"namespace": "", "time_ms": 1728606570227, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728606622252, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728606622252, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7098102495208356, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7098102495208356}}
+:::MLLOG {"namespace": "", "time_ms": 1728607773165, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728607826339, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728607826339, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7117401148361865, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7117401148361865}}
+:::MLLOG {"namespace": "", "time_ms": 1728608976638, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728609028698, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728609028698, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7135091588821822, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7135091588821822}}
+:::MLLOG {"namespace": "", "time_ms": 1728610184971, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728610238021, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728610238021, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.714804682570013, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.714804682570013}}
+:::MLLOG {"namespace": "", "time_ms": 1728611388386, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728611441326, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728611441326, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7155881319921319, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7155881319921319}}
+:::MLLOG {"namespace": "", "time_ms": 1728612596531, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728612648484, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728612648485, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7166956780410199, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7166956780410199}}
+:::MLLOG {"namespace": "", "time_ms": 1728613799458, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728613851376, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728613851376, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175573623673818, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7175573623673818}}
+:::MLLOG {"namespace": "", "time_ms": 1728615017853, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728615070888, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728615070888, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7177324852295528, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7177324852295528}}
+:::MLLOG {"namespace": "", "time_ms": 1728616224796, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728616276794, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728616276794, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7187149448529216, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7187149448529216}}
+:::MLLOG {"namespace": "", "time_ms": 1728617427166, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728617480097, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728617480098, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189440351084598, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7189440351084598}}
+:::MLLOG {"namespace": "", "time_ms": 1728618630518, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728618683523, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728618683524, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189651959015355, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7189651959015355}}
+:::MLLOG {"namespace": "", "time_ms": 1728619834438, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728619886484, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728619886484, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7197905807489396, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7197905807489396}}
+:::MLLOG {"namespace": "", "time_ms": 1728621047498, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728621100440, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728621100440, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7202326569216796, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7202326569216796}}
+:::MLLOG {"namespace": "", "time_ms": 1728621100441, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2999040}}
+:::MLLOG {"namespace": "", "time_ms": 1728621100441, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728621100441, "event_type": "POINT_IN_TIME", "key": "seed", "value": 18962, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,84 @@
+:::MLLOG {"namespace": "", "time_ms": 1728621113036, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113050, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113050, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113050, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113050, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113364, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728621113364, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728622708323, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728622718838, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734010, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734010, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734010, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734011, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734012, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728622734013, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728622779756, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728623989242, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728624047931, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728624047931, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38761837577562386, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38761837577562386}}
+:::MLLOG {"namespace": "", "time_ms": 1728625199877, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728625251862, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728625251862, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3991255365093096, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.3991255365093096}}
+:::MLLOG {"namespace": "", "time_ms": 1728626401253, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728626454238, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728626454239, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4629036359752662, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4629036359752662}}
+:::MLLOG {"namespace": "", "time_ms": 1728627603380, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728627655201, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728627655201, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5288661295558138, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5288661295558138}}
+:::MLLOG {"namespace": "", "time_ms": 1728628804682, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728628857572, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728628857573, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6336066501184932, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6336066501184932}}
+:::MLLOG {"namespace": "", "time_ms": 1728630005680, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728630057532, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728630057532, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.698150883303049, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.698150883303049}}
+:::MLLOG {"namespace": "", "time_ms": 1728631206492, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728631258444, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728631258444, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7087594178456637, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7087594178456637}}
+:::MLLOG {"namespace": "", "time_ms": 1728632406126, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728632458945, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728632458946, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.712074551420721, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.712074551420721}}
+:::MLLOG {"namespace": "", "time_ms": 1728633607305, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728633660325, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728633660325, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7133723366167564, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7133723366167564}}
+:::MLLOG {"namespace": "", "time_ms": 1728634807043, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728634859650, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728634859650, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7153185129380183, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7153185129380183}}
+:::MLLOG {"namespace": "", "time_ms": 1728636007670, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728636060454, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728636060454, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7153358636105497, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7153358636105497}}
+:::MLLOG {"namespace": "", "time_ms": 1728637212672, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728637266430, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728637266431, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7166292644147753, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7166292644147753}}
+:::MLLOG {"namespace": "", "time_ms": 1728638413489, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728638466302, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728638466303, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.717635804439778, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.717635804439778}}
+:::MLLOG {"namespace": "", "time_ms": 1728639621382, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728639673232, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728639673232, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7184983831576123, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7184983831576123}}
+:::MLLOG {"namespace": "", "time_ms": 1728640830900, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728640882936, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728640882936, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7185333441243461, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7185333441243461}}
+:::MLLOG {"namespace": "", "time_ms": 1728642038106, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728642090127, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728642090127, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.71883562362759, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.71883562362759}}
+:::MLLOG {"namespace": "", "time_ms": 1728643238601, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728643291575, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728643291575, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7195813601862262, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7195813601862262}}
+:::MLLOG {"namespace": "", "time_ms": 1728644439777, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728644491633, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728644491633, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7200058278167899, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7200058278167899}}
+:::MLLOG {"namespace": "", "time_ms": 1728644491633, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2699136}}
+:::MLLOG {"namespace": "", "time_ms": 1728644491633, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728644491633, "event_type": "POINT_IN_TIME", "key": "seed", "value": 7451, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,90 @@
+:::MLLOG {"namespace": "", "time_ms": 1728524927811, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927825, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927825, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927825, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927825, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927826, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728524927826, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728526542156, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728526553341, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570472, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570473, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570474, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570475, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728526570475, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728526620720, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728527836136, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728527898691, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728527898691, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3863042194326981, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.3863042194326981}}
+:::MLLOG {"namespace": "", "time_ms": 1728529058617, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728529115079, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728529115080, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.40961273735915393, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.40961273735915393}}
+:::MLLOG {"namespace": "", "time_ms": 1728530272495, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728530328049, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728530328049, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.43488811103970115, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.43488811103970115}}
+:::MLLOG {"namespace": "", "time_ms": 1728531485648, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728531541201, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728531541201, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4996365338081218, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.4996365338081218}}
+:::MLLOG {"namespace": "", "time_ms": 1728532697882, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728532753299, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728532753299, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5947849176998402, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5947849176998402}}
+:::MLLOG {"namespace": "", "time_ms": 1728533910552, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728533966102, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728533966102, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.690821249612306, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.690821249612306}}
+:::MLLOG {"namespace": "", "time_ms": 1728535121207, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728535177526, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728535177526, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7054878836201564, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7054878836201564}}
+:::MLLOG {"namespace": "", "time_ms": 1728536332776, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728536389219, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728536389220, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7097416914145819, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7097416914145819}}
+:::MLLOG {"namespace": "", "time_ms": 1728537544044, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728537599323, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728537599324, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7116119135763378, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7116119135763378}}
+:::MLLOG {"namespace": "", "time_ms": 1728538754453, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728538809839, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728538809839, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7137712529124653, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7137712529124653}}
+:::MLLOG {"namespace": "", "time_ms": 1728539966085, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728540021340, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728540021340, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7149044781941172, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7149044781941172}}
+:::MLLOG {"namespace": "", "time_ms": 1728541177404, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728541233606, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728541233606, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7153092085349753, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7153092085349753}}
+:::MLLOG {"namespace": "", "time_ms": 1728542397120, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728542452484, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728542452485, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.716667587340915, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.716667587340915}}
+:::MLLOG {"namespace": "", "time_ms": 1728543613320, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728543669699, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728543669699, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7173634540126506, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7173634540126506}}
+:::MLLOG {"namespace": "", "time_ms": 1728544826341, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728544881826, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728544881826, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7179233225148527, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7179233225148527}}
+:::MLLOG {"namespace": "", "time_ms": 1728546038641, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728546093965, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728546093965, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7184311815844229, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7184311815844229}}
+:::MLLOG {"namespace": "", "time_ms": 1728547260318, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728547316564, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728547316565, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7191096605050328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7191096605050328}}
+:::MLLOG {"namespace": "", "time_ms": 1728548476862, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728548532277, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728548532277, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192467427854418, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7192467427854418}}
+:::MLLOG {"namespace": "", "time_ms": 1728549694358, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728549750718, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728549750718, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7199938913341333, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7199938913341333}}
+:::MLLOG {"namespace": "", "time_ms": 1728550912201, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728550968747, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728550968748, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7203293324374981, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7203293324374981}}
+:::MLLOG {"namespace": "", "time_ms": 1728550968748, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2999040}}
+:::MLLOG {"namespace": "", "time_ms": 1728550968748, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728550968748, "event_type": "POINT_IN_TIME", "key": "seed", "value": 31643, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,96 @@
+:::MLLOG {"namespace": "", "time_ms": 1728550982004, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982018, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982018, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982018, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982018, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982200, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728550982201, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728552594185, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728552604676, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619008, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619008, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619008, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619009, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619010, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619010, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619010, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619011, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619011, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619011, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728552619011, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728552666986, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728553875969, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728553939951, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728553939951, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38616175677830683, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38616175677830683}}
+:::MLLOG {"namespace": "", "time_ms": 1728555092364, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728555150488, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728555150489, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.41162710727584095, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.41162710727584095}}
+:::MLLOG {"namespace": "", "time_ms": 1728556296541, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728556355574, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728556355575, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4350904176781545, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4350904176781545}}
+:::MLLOG {"namespace": "", "time_ms": 1728557502466, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728557561747, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728557561747, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.496310165579904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.496310165579904}}
+:::MLLOG {"namespace": "", "time_ms": 1728558708887, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728558766886, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728558766887, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5802412135699729, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5802412135699729}}
+:::MLLOG {"namespace": "", "time_ms": 1728559914092, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728559973260, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728559973260, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6652080704154694, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6652080704154694}}
+:::MLLOG {"namespace": "", "time_ms": 1728561120077, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728561179197, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728561179198, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6984256211339748, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.6984256211339748}}
+:::MLLOG {"namespace": "", "time_ms": 1728562325103, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728562383110, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728562383111, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7061839626088569, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7061839626088569}}
+:::MLLOG {"namespace": "", "time_ms": 1728563529512, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728563588342, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728563588343, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.709852535744663, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.709852535744663}}
+:::MLLOG {"namespace": "", "time_ms": 1728564732742, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728564791651, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728564791651, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7124425528931918, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7124425528931918}}
+:::MLLOG {"namespace": "", "time_ms": 1728565936303, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728565995377, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728565995377, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7139212561175718, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7139212561175718}}
+:::MLLOG {"namespace": "", "time_ms": 1728567148853, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728567207007, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728567207007, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7147917833311085, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7147917833311085}}
+:::MLLOG {"namespace": "", "time_ms": 1728568352291, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728568410261, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728568410261, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7158374547529306, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7158374547529306}}
+:::MLLOG {"namespace": "", "time_ms": 1728569560993, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728569619018, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728569619018, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7164395152271045, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7164395152271045}}
+:::MLLOG {"namespace": "", "time_ms": 1728570770470, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728570828476, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728570828476, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7171282240496328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7171282240496328}}
+:::MLLOG {"namespace": "", "time_ms": 1728571971989, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728572031038, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728572031039, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7176245095920047, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7176245095920047}}
+:::MLLOG {"namespace": "", "time_ms": 1728573181122, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728573239069, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728573239069, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7185501086189852, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7185501086189852}}
+:::MLLOG {"namespace": "", "time_ms": 1728574384118, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728574443334, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728574443334, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.718467390029532, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.718467390029532}}
+:::MLLOG {"namespace": "", "time_ms": 1728575608940, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728575667042, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728575667042, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7193740987105504, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7193740987105504}}
+:::MLLOG {"namespace": "", "time_ms": 1728576812654, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728576871686, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728576871686, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.719547293968521, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.719547293968521}}
+:::MLLOG {"namespace": "", "time_ms": 1728578016561, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3148992, "step_num": 47712}}
+:::MLLOG {"namespace": "", "time_ms": 1728578075504, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3148992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3148992, "step_num": 47712, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728578075504, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198685481652716, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3148992, "masked_lm_accuracy": 0.7198685481652716}}
+:::MLLOG {"namespace": "", "time_ms": 1728579221384, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 3298944, "step_num": 49984}}
+:::MLLOG {"namespace": "", "time_ms": 1728579280495, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 3298944, "step_num": 49984, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728579280495, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7202844063393284, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 3298944, "masked_lm_accuracy": 0.7202844063393284}}
+:::MLLOG {"namespace": "", "time_ms": 1728579280496, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 3298944, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 3298944}}
+:::MLLOG {"namespace": "", "time_ms": 1728579280496, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728579280496, "event_type": "POINT_IN_TIME", "key": "seed", "value": 26715, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,87 @@
+:::MLLOG {"namespace": "", "time_ms": 1728579292347, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292360, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292360, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292360, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292360, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292510, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728579292511, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728580879014, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728580889187, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903804, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903804, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903805, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728580903806, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728580951866, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728582174949, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728582235954, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728582235955, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38839849547204247, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38839849547204247}}
+:::MLLOG {"namespace": "", "time_ms": 1728583404398, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728583457118, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728583457118, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4043779520660943, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4043779520660943}}
+:::MLLOG {"namespace": "", "time_ms": 1728584623300, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728584677358, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728584677358, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4277538998833515, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.4277538998833515}}
+:::MLLOG {"namespace": "", "time_ms": 1728585841710, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728585894585, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728585894586, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4895906483404781, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.4895906483404781}}
+:::MLLOG {"namespace": "", "time_ms": 1728587058762, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728587112509, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728587112509, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5868821954505488, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.5868821954505488}}
+:::MLLOG {"namespace": "", "time_ms": 1728588276137, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728588329834, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728588329834, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6855756070131875, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6855756070131875}}
+:::MLLOG {"namespace": "", "time_ms": 1728589494361, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728589548276, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728589548276, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7039273525328428, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7039273525328428}}
+:::MLLOG {"namespace": "", "time_ms": 1728590710884, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728590765567, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728590765567, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7093820230075536, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7093820230075536}}
+:::MLLOG {"namespace": "", "time_ms": 1728591929060, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728591981842, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728591981842, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7118489585669369, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7118489585669369}}
+:::MLLOG {"namespace": "", "time_ms": 1728593146668, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728593199377, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728593199377, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7138080239009914, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7138080239009914}}
+:::MLLOG {"namespace": "", "time_ms": 1728594362906, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728594416691, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728594416692, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7149828011287352, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7149828011287352}}
+:::MLLOG {"namespace": "", "time_ms": 1728595585602, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728595639450, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728595639450, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7161310530023511, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7161310530023511}}
+:::MLLOG {"namespace": "", "time_ms": 1728596808968, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728596861749, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728596861749, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7167342547606621, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7167342547606621}}
+:::MLLOG {"namespace": "", "time_ms": 1728598025777, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728598079566, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728598079567, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175075857669347, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7175075857669347}}
+:::MLLOG {"namespace": "", "time_ms": 1728599242962, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728599297649, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728599297649, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7177893112502416, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7177893112502416}}
+:::MLLOG {"namespace": "", "time_ms": 1728600461272, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728600514019, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728600514019, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7186980957604484, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7186980957604484}}
+:::MLLOG {"namespace": "", "time_ms": 1728601688118, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728601741703, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728601741703, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7194635579882086, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7194635579882086}}
+:::MLLOG {"namespace": "", "time_ms": 1728602912396, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728602965140, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728602965140, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7196329377980453, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7196329377980453}}
+:::MLLOG {"namespace": "", "time_ms": 1728604140305, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728604193177, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728604193177, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7201142786646146, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7201142786646146}}
+:::MLLOG {"namespace": "", "time_ms": 1728604193178, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2849088}}
+:::MLLOG {"namespace": "", "time_ms": 1728604193178, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728604193178, "event_type": "POINT_IN_TIME", "key": "seed", "value": 24659, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,90 @@
+:::MLLOG {"namespace": "", "time_ms": 1728604206238, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206251, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206251, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206251, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206251, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206399, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728604206400, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728605831206, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728605841697, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856169, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856169, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856170, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728605856171, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728605907083, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728607116687, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728607176173, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728607176173, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38733550816792245, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38733550816792245}}
+:::MLLOG {"namespace": "", "time_ms": 1728608331645, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728608385445, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728608385445, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.4099236127138853, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.4099236127138853}}
+:::MLLOG {"namespace": "", "time_ms": 1728609536865, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728609589476, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728609589476, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.46445941542463526, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.46445941542463526}}
+:::MLLOG {"namespace": "", "time_ms": 1728610742733, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728610796400, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728610796401, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5356849022231992, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.5356849022231992}}
+:::MLLOG {"namespace": "", "time_ms": 1728611950463, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728612003989, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728612003989, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6243599510936588, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.6243599510936588}}
+:::MLLOG {"namespace": "", "time_ms": 1728613157626, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728613211058, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728613211058, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6915194793740074, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.6915194793740074}}
+:::MLLOG {"namespace": "", "time_ms": 1728614363023, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728614416397, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728614416397, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.70565997331387, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.70565997331387}}
+:::MLLOG {"namespace": "", "time_ms": 1728615568690, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728615622127, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728615622128, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7094783287385873, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7094783287385873}}
+:::MLLOG {"namespace": "", "time_ms": 1728616773489, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728616825919, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728616825919, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7119873171805191, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7119873171805191}}
+:::MLLOG {"namespace": "", "time_ms": 1728617978563, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728618032221, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728618032222, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7142333673658525, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7142333673658525}}
+:::MLLOG {"namespace": "", "time_ms": 1728619183815, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728619236305, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728619236305, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7149768588638763, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7149768588638763}}
+:::MLLOG {"namespace": "", "time_ms": 1728620386747, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728620439998, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728620439998, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7156523679928931, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7156523679928931}}
+:::MLLOG {"namespace": "", "time_ms": 1728621597193, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728621650622, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728621650622, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7164913090174971, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7164913090174971}}
+:::MLLOG {"namespace": "", "time_ms": 1728622809144, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728622861570, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728622861570, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7172874884828523, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7172874884828523}}
+:::MLLOG {"namespace": "", "time_ms": 1728624015271, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728624067817, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728624067818, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7180004538213985, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7180004538213985}}
+:::MLLOG {"namespace": "", "time_ms": 1728625221988, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728625275277, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728625275277, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7182580689577264, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7182580689577264}}
+:::MLLOG {"namespace": "", "time_ms": 1728626428374, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2549184, "step_num": 38624}}
+:::MLLOG {"namespace": "", "time_ms": 1728626481696, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2549184, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2549184, "step_num": 38624, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728626481696, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7189949285361891, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2549184, "masked_lm_accuracy": 0.7189949285361891}}
+:::MLLOG {"namespace": "", "time_ms": 1728627644056, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2699136, "step_num": 40896}}
+:::MLLOG {"namespace": "", "time_ms": 1728627696455, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2699136, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2699136, "step_num": 40896, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728627696456, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7192161206220823, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2699136, "masked_lm_accuracy": 0.7192161206220823}}
+:::MLLOG {"namespace": "", "time_ms": 1728628864498, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2849088, "step_num": 43168}}
+:::MLLOG {"namespace": "", "time_ms": 1728628917907, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2849088, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2849088, "step_num": 43168, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728628917908, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7198127043697744, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2849088, "masked_lm_accuracy": 0.7198127043697744}}
+:::MLLOG {"namespace": "", "time_ms": 1728630073687, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2999040, "step_num": 45440}}
+:::MLLOG {"namespace": "", "time_ms": 1728630126070, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2999040, "step_num": 45440, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728630126070, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7204812415431342, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2999040, "masked_lm_accuracy": 0.7204812415431342}}
+:::MLLOG {"namespace": "", "time_ms": 1728630126070, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2999040, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2999040}}
+:::MLLOG {"namespace": "", "time_ms": 1728630126071, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728630126071, "event_type": "POINT_IN_TIME", "key": "seed", "value": 6018, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -0,0 +1,78 @@
+:::MLLOG {"namespace": "", "time_ms": 1728630138587, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 631}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138600, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox_red", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 632}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138600, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 633}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138601, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 634}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138601, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "bert", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 636}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138750, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 639}}
+:::MLLOG {"namespace": "", "time_ms": 1728630138750, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 640}}
+:::MLLOG {"namespace": "", "time_ms": 1728631762258, "event_type": "POINT_IN_TIME", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 842}}
+:::MLLOG {"namespace": "", "time_ms": 1728631772243, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 643}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787279, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 66, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 711}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 512, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 712}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "max_predictions_per_seq", "value": 76, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 713}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "LAMB", "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 715}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 716}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "opt_lamb_weight_decay_rate", "value": 0.01, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 717}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787280, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_1", "value": 0.9, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 718}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "opt_lamb_beta_2", "value": 0.999, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 719}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "opt_lamb_learning_rate_decay_poly_power", "value": 1.0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 720}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "opt_lamb_epsilon", "value": 1e-06, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 721}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 723}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 724}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "start_warmup_step", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 725}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787281, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_training_steps", "value": 55000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 726}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787282, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 1, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 727}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787282, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 10002, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 728}}
+:::MLLOG {"namespace": "", "time_ms": 1728631787282, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 3630000, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 729}}
+:::MLLOG {"namespace": "", "time_ms": 1728631833698, "event_type": "INTERVAL_START", "key": "epoch_start", "value": 0, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 753, "epoch_num": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1728633040032, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 149952, "step_num": 2272}}
+:::MLLOG {"namespace": "", "time_ms": 1728633099704, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 149952, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 149952, "step_num": 2272, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728633099705, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.38797098084452913, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 149952, "masked_lm_accuracy": 0.38797098084452913}}
+:::MLLOG {"namespace": "", "time_ms": 1728634252081, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 299904, "step_num": 4544}}
+:::MLLOG {"namespace": "", "time_ms": 1728634305117, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 299904, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 299904, "step_num": 4544, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728634305117, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.3995994811617263, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 299904, "masked_lm_accuracy": 0.3995994811617263}}
+:::MLLOG {"namespace": "", "time_ms": 1728635455137, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 449856, "step_num": 6816}}
+:::MLLOG {"namespace": "", "time_ms": 1728635508121, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 449856, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 449856, "step_num": 6816, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728635508121, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.5011176564375941, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 449856, "masked_lm_accuracy": 0.5011176564375941}}
+:::MLLOG {"namespace": "", "time_ms": 1728636656616, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 599808, "step_num": 9088}}
+:::MLLOG {"namespace": "", "time_ms": 1728636710601, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 599808, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 599808, "step_num": 9088, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728636710602, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.6374410459147146, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 599808, "masked_lm_accuracy": 0.6374410459147146}}
+:::MLLOG {"namespace": "", "time_ms": 1728637859253, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 749760, "step_num": 11360}}
+:::MLLOG {"namespace": "", "time_ms": 1728637913116, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 749760, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 749760, "step_num": 11360, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728637913117, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7006531824018688, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 749760, "masked_lm_accuracy": 0.7006531824018688}}
+:::MLLOG {"namespace": "", "time_ms": 1728639062400, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 899712, "step_num": 13632}}
+:::MLLOG {"namespace": "", "time_ms": 1728639116398, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 899712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 899712, "step_num": 13632, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728639116398, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7092244053358937, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 899712, "masked_lm_accuracy": 0.7092244053358937}}
+:::MLLOG {"namespace": "", "time_ms": 1728640264599, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1049664, "step_num": 15904}}
+:::MLLOG {"namespace": "", "time_ms": 1728640317583, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1049664, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1049664, "step_num": 15904, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728640317584, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7131378053546167, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1049664, "masked_lm_accuracy": 0.7131378053546167}}
+:::MLLOG {"namespace": "", "time_ms": 1728641465346, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1199616, "step_num": 18176}}
+:::MLLOG {"namespace": "", "time_ms": 1728641518225, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1199616, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1199616, "step_num": 18176, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728641518225, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7143681931080901, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1199616, "masked_lm_accuracy": 0.7143681931080901}}
+:::MLLOG {"namespace": "", "time_ms": 1728642666024, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1349568, "step_num": 20448}}
+:::MLLOG {"namespace": "", "time_ms": 1728642718905, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1349568, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1349568, "step_num": 20448, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728642718905, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7157575550305322, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1349568, "masked_lm_accuracy": 0.7157575550305322}}
+:::MLLOG {"namespace": "", "time_ms": 1728643864313, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1499520, "step_num": 22720}}
+:::MLLOG {"namespace": "", "time_ms": 1728643917070, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1499520, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1499520, "step_num": 22720, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728643917070, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7172654897564532, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1499520, "masked_lm_accuracy": 0.7172654897564532}}
+:::MLLOG {"namespace": "", "time_ms": 1728645063069, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1649472, "step_num": 24992}}
+:::MLLOG {"namespace": "", "time_ms": 1728645115906, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1649472, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1649472, "step_num": 24992, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728645115907, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7175817649094158, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1649472, "masked_lm_accuracy": 0.7175817649094158}}
+:::MLLOG {"namespace": "", "time_ms": 1728646261029, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1799424, "step_num": 27264}}
+:::MLLOG {"namespace": "", "time_ms": 1728646314945, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1799424, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1799424, "step_num": 27264, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728646314946, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7180203909851078, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1799424, "masked_lm_accuracy": 0.7180203909851078}}
+:::MLLOG {"namespace": "", "time_ms": 1728647468431, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 1949376, "step_num": 29536}}
+:::MLLOG {"namespace": "", "time_ms": 1728647522632, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 1949376, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 1949376, "step_num": 29536, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728647522633, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7191875515115712, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 1949376, "masked_lm_accuracy": 0.7191875515115712}}
+:::MLLOG {"namespace": "", "time_ms": 1728648686121, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2099328, "step_num": 31808}}
+:::MLLOG {"namespace": "", "time_ms": 1728648739069, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2099328, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2099328, "step_num": 31808, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728648739069, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7196067409309428, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2099328, "masked_lm_accuracy": 0.7196067409309428}}
+:::MLLOG {"namespace": "", "time_ms": 1728649885357, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2249280, "step_num": 34080}}
+:::MLLOG {"namespace": "", "time_ms": 1728649939244, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2249280, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2249280, "step_num": 34080, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728649939245, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7197489482454004, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2249280, "masked_lm_accuracy": 0.7197489482454004}}
+:::MLLOG {"namespace": "", "time_ms": 1728651085252, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 806, "epoch_num": 2399232, "step_num": 36352}}
+:::MLLOG {"namespace": "", "time_ms": 1728651139304, "event_type": "INTERVAL_END", "key": "eval_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 861, "epoch_count": 2399232, "step_num": 36352, "samples_count": 10002}}
+:::MLLOG {"namespace": "", "time_ms": 1728651139304, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 0.7202936197442785, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 862, "epoch_num": 2399232, "masked_lm_accuracy": 0.7202936197442785}}
+:::MLLOG {"namespace": "", "time_ms": 1728651139305, "event_type": "POINT_IN_TIME", "key": "epoch_stop", "value": 2399232, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 880, "epoch_num": 2399232}}
+:::MLLOG {"namespace": "", "time_ms": 1728651139305, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 881, "status": "success"}}
+:::MLLOG {"namespace": "", "time_ms": 1728651139305, "event_type": "POINT_IN_TIME", "key": "seed", "value": 32021, "metadata": {"file": "tinygrad/examples/mlperf/model_train.py", "lineno": 882}}
@@ -1,12 +1,13 @@
 {
  "submitter": "tinycorp",
  "division": "closed",
-  "status": "available",
+  "status": "Available on-premise",
  "system_name": "tinybox green",
  "number_of_nodes": "1",
  "host_processors_per_node": "1",
  "host_processor_model_name": "AMD EPYC 7532 32-Core Processor",
-  "host_processor_core_count": "64",
+  "host_processor_core_count": "32",
+  "host_processor_vcpu_count": "64",
  "host_processor_frequency": "",
  "host_processor_caches": "",
  "host_processor_interconnect": "",
@@ -27,7 +28,7 @@
  "accelerator_interconnect_topology": "",
  "cooling": "air",
  "hw_notes": "",
-  "framework": "tinygrad, commit 0e8aa0e2886bf9a2d3ce093bce87305e182e6d4a",
+  "framework": "tinygrad, commit b5546912e24e0a864b35924da4efa5d71cfe368b",
  "other_software_stack": {
    "python": "3.10.12",
    "CUDA": "12.4"
@@ -1,12 +1,13 @@
 {
  "submitter": "tinycorp",
  "division": "closed",
-  "status": "available",
+  "status": "Available on-premise",
  "system_name": "tinybox red",
  "number_of_nodes": "1",
  "host_processors_per_node": "1",
  "host_processor_model_name": "AMD EPYC 7532 32-Core Processor",
-  "host_processor_core_count": "64",
+  "host_processor_core_count": "32",
+  "host_processor_vcpu_count": "64",
  "host_processor_frequency": "",
  "host_processor_caches": "",
  "host_processor_interconnect": "",
@@ -27,10 +28,10 @@
  "accelerator_interconnect_topology": "",
  "cooling": "air",
  "hw_notes": "",
-  "framework": "tinygrad, commit 0e8aa0e2886bf9a2d3ce093bce87305e182e6d4a",
+  "framework": "tinygrad, commit b5546912e24e0a864b35924da4efa5d71cfe368b",
  "other_software_stack": {
    "python": "3.10.12",
-    "ROCm": "6.1"
+    "ROCm": "6.1.3"
  },
  "operating_system": "Ubuntu 22.04.4",
  "sw_notes": ""
@@ -18,7 +18,7 @@ from tinygrad.device import Buffer
 from tinygrad.helpers import partition, Context, fetch, getenv, DEBUG, tqdm
 from tinygrad.engine.realize import run_schedule, lower_schedule, ExecItem, CompiledRunner, memory_planner
 from tinygrad.engine.schedule import ScheduleItem, create_schedule
-from tinygrad.ops import MetaOps
+from tinygrad.ops import UOps
 from tinygrad.tensor import _to_np_dtype
 Device.DEFAULT = "GPU"

@@ -49,7 +49,7 @@ def get_schedule(onnx_data) -> Tuple[List[ScheduleItem], List[ScheduleItem]]:
  print(f"{len(schedule)} schedule items depend on the input, {len(schedule_independent)} don't")

  # confirm no non-sink metaop in the (non independent) schedule except for the ones that load the input buffers
-  assert all(si.ast.op is MetaOps.KERNEL or out in input_lb for si in schedule for out in si.outputs), "has non SINK ops, can't compile to Thneed"
+  assert all(si.ast.op is UOps.SINK or out in input_lb for si in schedule for out in si.outputs), "has non SINK ops, can't compile to Thneed"
  return schedule, schedule_independent, inputs

 def test_vs_onnx(onnx_data, eis:Optional[List[ExecItem]], inputs:Dict[str, Tensor]):
@@ -105,7 +105,7 @@ if __name__ == "__main__":
  #exit(0)

  schedule, schedule_independent, inputs = get_schedule(onnx_data)
-  schedule, schedule_input = partition(schedule, lambda x: x.ast.op is MetaOps.KERNEL)
+  schedule, schedule_input = partition(schedule, lambda x: x.ast.op is UOps.SINK)
  print(f"{len(schedule_input)} inputs")

  run_schedule(schedule_independent)
@@ -164,6 +164,7 @@ if __name__ == "__main__":

  saved_binaries = set()
  binaries = []
+  gated_read_image_count = 0
  GlobalCounters.reset()
  with Context(DEBUG=max(DEBUG.value, 2)):
    for ei in eis:
@@ -173,6 +174,7 @@ if __name__ == "__main__":
        jdat['binaries'].append({"name":prg.p.function_name, "length":len(prg.lib)})
        binaries.append(prg.lib)
        saved_binaries.add(prg.p.function_name)
+      gated_read_image_count += prg.p.src.count("?read_image")
      ei.run()
      jdat['kernels'].append({
        "name": prg.p.function_name,
@@ -184,6 +186,10 @@ if __name__ == "__main__":
        "arg_size": [8]*len(ei.bufs),
      })

+  if (allowed_gated_read_image:=getenv("ALLOWED_GATED_READ_IMAGE", -1)) != -1:
+    assert gated_read_image_count <= allowed_gated_read_image, \
+      f"too many gated read_image! {gated_read_image_count=}, {allowed_gated_read_image=}"
+
  output_fn = sys.argv[2] if len(sys.argv) >= 3 else "/tmp/output.thneed"
  print(f"saving thneed to {output_fn} with {len(weights)} buffers and {len(binaries)} binaries")
  with open(output_fn, "wb") as f:
@@ -0,0 +1,80 @@
+import os, sys, pickle, time
+import numpy as np
+if "FLOAT16" not in os.environ: os.environ["FLOAT16"] = "1"
+if "IMAGE" not in os.environ: os.environ["IMAGE"] = "2"
+if "NOLOCALS" not in os.environ: os.environ["NOLOCALS"] = "1"
+if "JIT_BATCH_SIZE" not in os.environ: os.environ["JIT_BATCH_SIZE"] = "0"
+
+from tinygrad import fetch, Tensor, TinyJit, Device, Context, GlobalCounters
+from tinygrad.helpers import OSX, DEBUG, getenv
+from tinygrad.tensor import _from_np_dtype
+
+import onnx
+from onnx.helper import tensor_dtype_to_np_dtype
+from extra.onnx import get_run_onnx   # TODO: port to main tinygrad
+
+OPENPILOT_MODEL = sys.argv[1] if len(sys.argv) > 1 else "https://github.com/commaai/openpilot/raw/v0.9.7/selfdrive/modeld/models/supercombo.onnx"
+OUTPUT = "/tmp/openpilot.pkl"
+
+def compile():
+  # hack to fix GPU on OSX: max doesn't work on half, see test/external/external_gpu_fail_osx.py
+  if OSX:
+    from tinygrad.ops import BinaryOps
+    from tinygrad.renderer.cstyle import ClangRenderer, CStyleLanguage
+    CStyleLanguage.code_for_op[BinaryOps.MAX] = ClangRenderer.code_for_op[BinaryOps.MAX]
+
+  Tensor.no_grad = True
+  Tensor.training = False
+
+  onnx_bytes = fetch(OPENPILOT_MODEL)
+  onnx_model = onnx.load(onnx_bytes)
+  run_onnx = get_run_onnx(onnx_model)
+  print("loaded model")
+
+  input_shapes = {inp.name:tuple(x.dim_value for x in inp.type.tensor_type.shape.dim) for inp in onnx_model.graph.input}
+  input_types = {inp.name: tensor_dtype_to_np_dtype(inp.type.tensor_type.elem_type) for inp in onnx_model.graph.input}
+  Tensor.manual_seed(100)
+  new_inputs = {k:Tensor.randn(*shp, dtype=_from_np_dtype(input_types[k])).mul(8).realize() for k,shp in sorted(input_shapes.items())}
+  print("created tensors")
+
+  run_onnx_jit = TinyJit(lambda **kwargs: run_onnx(kwargs), prune=True)
+  for i in range(3):
+    GlobalCounters.reset()
+    print(f"run {i}")
+    with Context(DEBUG=max(DEBUG.value, 2 if i == 2 else 1)):
+      ret = next(iter(run_onnx_jit(**new_inputs).values())).cast('float32').numpy()
+    if i == 0: test_val = np.copy(ret)
+  print(f"captured {len(run_onnx_jit.captured.jit_cache)} kernels")
+  np.testing.assert_equal(test_val, ret)
+  print("jit run validated")
+
+  with open(OUTPUT, "wb") as f:
+    pickle.dump(run_onnx_jit, f)
+  mdl_sz = os.path.getsize(onnx_bytes)
+  pkl_sz = os.path.getsize(OUTPUT)
+  print(f"mdl size is {mdl_sz/1e6:.2f}M")
+  print(f"pkl size is {pkl_sz/1e6:.2f}M")
+  print("**** compile done ****")
+  return test_val
+
+def test(test_val=None):
+  with open(OUTPUT, "rb") as f:
+    run = pickle.load(f)
+  Tensor.manual_seed(100)
+  new_inputs = {nm:Tensor.randn(*st.shape, dtype=dtype).mul(8).realize() for nm, (st, _, dtype, _) in
+                sorted(zip(run.captured.expected_names, run.captured.expected_st_vars_dtype_device))}
+  for _ in range(20):
+    st = time.perf_counter()
+    out = run(**new_inputs)
+    mt = time.perf_counter()
+    val = out['outputs'].numpy()
+    et = time.perf_counter()
+    print(f"enqueue {(mt-st)*1e3:6.2f} ms -- total run {(et-st)*1e3:6.2f} ms")
+  print(out, val.shape, val.dtype)
+  if test_val is not None: np.testing.assert_equal(test_val, val)
+  print("**** test done ****")
+
+if __name__ == "__main__":
+  test_val = compile() if not getenv("RUN") else None
+  test(test_val)
+
@@ -12,7 +12,7 @@ from extra.models.unet import UNetModel, Upsample, Downsample, timestep_embeddin
 from examples.stable_diffusion import ResnetBlock, Mid
 import numpy as np

-from typing import Dict, List, Callable, Optional, Any, Set, Tuple
+from typing import Dict, List, Callable, Optional, Any, Set, Tuple, Union, Type
 import argparse, tempfile
 from abc import ABC, abstractmethod
 from pathlib import Path
@@ -47,21 +47,14 @@ class DiffusionModel:
    self.diffusion_model = UNetModel(*args, **kwargs)


-class Embedder(ABC):
-  input_key: str
-  @abstractmethod
-  def __call__(self, x:Tensor) -> Tensor:
-    pass
-
-
 # https://github.com/Stability-AI/generative-models/blob/fbdc58cab9f4ee2be7a5e1f2e2787ecd9311942f/sgm/modules/encoders/modules.py#L913
 class ConcatTimestepEmbedderND(Embedder):
  def __init__(self, outdim:int, input_key:str):
    self.outdim = outdim
    self.input_key = input_key

-  def __call__(self, x:Tensor):
-    assert len(x.shape) == 2
+  def __call__(self, x:Union[str,List[str],Tensor]):
+    assert isinstance(x, Tensor) and len(x.shape) == 2
    emb = timestep_embedding(x.flatten(), self.outdim)
    emb = emb.reshape((x.shape[0],-1))
    return emb
@@ -91,9 +84,8 @@ class Conditioner:
      emb_out = embedder(batch[embedder.input_key])

      if isinstance(emb_out, Tensor):
-        emb_out = [emb_out]
-      else:
-        assert isinstance(emb_out, (list, tuple))
+        emb_out = (emb_out,)
+      assert isinstance(emb_out, (list, tuple))

      for emb in emb_out:
        if embedder.input_key in force_zero_embeddings:
@@ -248,22 +240,23 @@ class SDXL:
    self.sigmas = self.discretization(config["denoiser"]["num_idx"], flip=True)

  # https://github.com/Stability-AI/generative-models/blob/fbdc58cab9f4ee2be7a5e1f2e2787ecd9311942f/sgm/inference/helpers.py#L173
-  def create_conditioning(self, pos_prompt:str, img_width:int, img_height:int, aesthetic_score:float=5.0) -> Tuple[Dict,Dict]:
+  def create_conditioning(self, pos_prompts:List[str], img_width:int, img_height:int, aesthetic_score:float=5.0) -> Tuple[Dict,Dict]:
+    N = len(pos_prompts)
    batch_c : Dict = {
-      "txt": pos_prompt,
+      "txt": pos_prompts,
      "original_size_as_tuple": Tensor([img_height,img_width]).repeat(N,1),
      "crop_coords_top_left": Tensor([0,0]).repeat(N,1),
      "target_size_as_tuple": Tensor([img_height,img_width]).repeat(N,1),
      "aesthetic_score": Tensor([aesthetic_score]).repeat(N,1),
    }
    batch_uc: Dict = {
-      "txt": "",
+      "txt": [""]*N,
      "original_size_as_tuple": Tensor([img_height,img_width]).repeat(N,1),
      "crop_coords_top_left": Tensor([0,0]).repeat(N,1),
      "target_size_as_tuple": Tensor([img_height,img_width]).repeat(N,1),
      "aesthetic_score": Tensor([aesthetic_score]).repeat(N,1),
    }
-    return model.conditioner(batch_c), model.conditioner(batch_uc, force_zero_embeddings=["txt"])
+    return self.conditioner(batch_c), self.conditioner(batch_uc, force_zero_embeddings=["txt"])

  # https://github.com/Stability-AI/generative-models/blob/fbdc58cab9f4ee2be7a5e1f2e2787ecd9311942f/sgm/modules/diffusionmodules/denoiser.py#L42
  def denoise(self, x:Tensor, sigma:Tensor, cond:Dict) -> Tensor:
@@ -289,19 +282,29 @@ class SDXL:
    return self.first_stage_model.decode(1.0 / 0.13025 * x)


-class VanillaCFG:
+class Guider(ABC):
  def __init__(self, scale:float):
    self.scale = scale

-  def prepare_inputs(self, x:Tensor, s:float, c:Dict, uc:Dict) -> Tuple[Tensor,Tensor,Tensor]:
+  @abstractmethod
+  def __call__(self, denoiser, x:Tensor, s:Tensor, c:Dict, uc:Dict) -> Tensor:
+    pass
+
+class VanillaCFG(Guider):
+  def __call__(self, denoiser, x:Tensor, s:Tensor, c:Dict, uc:Dict) -> Tensor:
    c_out = {}
    for k in c:
      assert k in ["vector", "crossattn", "concat"]
      c_out[k] = Tensor.cat(uc[k], c[k], dim=0)
-    return Tensor.cat(x, x), Tensor.cat(s, s), c_out

-  def __call__(self, x:Tensor, sigma:float) -> Tensor:
-    x_u, x_c = x.chunk(2)
+    x_u, x_c = denoiser(Tensor.cat(x, x), Tensor.cat(s, s), c_out).chunk(2)
+    x_pred = x_u + self.scale*(x_c - x_u)
+    return x_pred
+
+class SplitVanillaCFG(Guider):
+  def __call__(self, denoiser, x:Tensor, s:Tensor, c:Dict, uc:Dict) -> Tensor:
+    x_u = denoiser(x, s, uc)
+    x_c = denoiser(x, s, c)
    x_pred = x_u + self.scale*(x_c - x_u)
    return x_pred

@@ -309,13 +312,12 @@ class VanillaCFG:
 # https://github.com/Stability-AI/generative-models/blob/fbdc58cab9f4ee2be7a5e1f2e2787ecd9311942f/sgm/modules/diffusionmodules/sampling.py#L21
 # https://github.com/Stability-AI/generative-models/blob/fbdc58cab9f4ee2be7a5e1f2e2787ecd9311942f/sgm/modules/diffusionmodules/sampling.py#L287
 class DPMPP2MSampler:
-  def __init__(self, cfg_scale:float):
+  def __init__(self, cfg_scale:float, guider_cls:Type[Guider]=VanillaCFG):
    self.discretization = LegacyDDPMDiscretization()
-    self.guider = VanillaCFG(cfg_scale)
+    self.guider = guider_cls(cfg_scale)

  def sampler_step(self, old_denoised:Optional[Tensor], prev_sigma:Optional[Tensor], sigma:Tensor, next_sigma:Tensor, denoiser, x:Tensor, c:Dict, uc:Dict) -> Tuple[Tensor,Tensor]:
-    denoised = denoiser(*self.guider.prepare_inputs(x, sigma, c, uc))
-    denoised = self.guider(denoised, sigma)
+    denoised = self.guider(denoiser, x, sigma, c, uc)

    t, t_next = sigma.log().neg(), next_sigma.log().neg()
    h = t_next - t
@@ -336,7 +338,7 @@ class DPMPP2MSampler:
    return x, denoised

  def __call__(self, denoiser, x:Tensor, c:Dict, uc:Dict, num_steps:int, timing=False) -> Tensor:
-    sigmas = self.discretization(num_steps)
+    sigmas = self.discretization(num_steps).to(x.device)
    x *= Tensor.sqrt(1.0 + sigmas[0] ** 2.0)
    num_sigmas = len(sigmas)

@@ -345,9 +347,9 @@ class DPMPP2MSampler:
      with Timing("step in ", enabled=timing, on_exit=lambda _: f", using {GlobalCounters.mem_used/1e9:.2f} GB"):
        x, old_denoised = self.sampler_step(
          old_denoised=old_denoised,
-          prev_sigma=(None if i==0 else sigmas[i-1].reshape(x.shape[0])),
-          sigma=sigmas[i].reshape(x.shape[0]),
-          next_sigma=sigmas[i+1].reshape(x.shape[0]),
+          prev_sigma=(None if i==0 else sigmas[i-1].expand(x.shape[0])),
+          sigma=sigmas[i].expand(x.shape[0]),
+          next_sigma=sigmas[i+1].expand(x.shape[0]),
          denoiser=denoiser,
          x=x,
          c=c,
@@ -391,7 +393,7 @@ if __name__ == "__main__":
  assert args.width  % F == 0, f"img_width must be multiple of {F}, got {args.width}"
  assert args.height % F == 0, f"img_height must be multiple of {F}, got {args.height}"

-  c, uc = model.create_conditioning(args.prompt, args.width, args.height)
+  c, uc = model.create_conditioning([args.prompt], args.width, args.height)
  del model.conditioner
  for v in c .values(): v.realize()
  for v in uc.values(): v.realize()
@@ -423,6 +425,6 @@ if __name__ == "__main__":
  if args.prompt == default_prompt and args.steps == 10 and args.seed == 0 and args.guidance == 6.0 and args.width == args.height == 1024 \
    and not args.weights:
    ref_image = Tensor(np.array(Image.open(Path(__file__).parent / "sdxl_seed0.png")))
-    distance = (((x - ref_image).cast(dtypes.float) / ref_image.max())**2).mean().item()
+    distance = (((x.cast(dtypes.float) - ref_image.cast(dtypes.float)) / ref_image.max())**2).mean().item()
    assert distance < 4e-3, colored(f"validation failed with {distance=}", "red")
    print(colored(f"output validated with {distance=}", "green"))
@@ -218,7 +218,7 @@ class StableDiffusion:
 if __name__ == "__main__":
  default_prompt = "a horse sized cat eating a bagel"
  parser = argparse.ArgumentParser(description='Run Stable Diffusion', formatter_class=argparse.ArgumentDefaultsHelpFormatter)
-  parser.add_argument('--steps', type=int, default=5, help="Number of steps in diffusion")
+  parser.add_argument('--steps', type=int, default=6, help="Number of steps in diffusion")
  parser.add_argument('--prompt', type=str, default=default_prompt, help="Phrase to render")
  parser.add_argument('--out', type=str, default=Path(tempfile.gettempdir()) / "rendered.png", help="Output filename")
  parser.add_argument('--noshow', action='store_true', help="Don't show the image")
@@ -287,8 +287,8 @@ if __name__ == "__main__":
  if not args.noshow: im.show()

  # validation!
-  if args.prompt == default_prompt and args.steps == 5 and args.seed == 0 and args.guidance == 7.5:
+  if args.prompt == default_prompt and args.steps == 6 and args.seed == 0 and args.guidance == 7.5:
    ref_image = Tensor(np.array(Image.open(Path(__file__).parent / "stable_diffusion_seed0.png")))
-    distance = (((x - ref_image).cast(dtypes.float) / ref_image.max())**2).mean().item()
-    assert distance < 3e-4, colored(f"validation failed with {distance=}", "red")
+    distance = (((x.cast(dtypes.float) - ref_image.cast(dtypes.float)) / ref_image.max())**2).mean().item()
+    assert distance < 50e-5, colored(f"validation failed with {distance=}", "red")
    print(colored(f"output validated with {distance=}", "green"))
@@ -0,0 +1,53 @@
+# beautiful mnist in the new "one-shot" style
+# one realize in the whole graph
+# depends on:
+#  - "big graph" UOp scheduling
+#  - symbolic removal
+
+from examples.beautiful_mnist import Model
+from tinygrad import Tensor, nn, getenv, GlobalCounters
+from tinygrad.nn.datasets import mnist
+from tinygrad.helpers import trange, DEBUG
+
+if __name__ == "__main__":
+  X_train, Y_train, X_test, Y_test = mnist()
+  print("*** got data")
+
+  model = Model()
+  print("*** got model")
+
+  opt = nn.optim.Adam(nn.state.get_parameters(model))
+  print("*** got optimizer")
+
+  samples = Tensor.randint(getenv("STEPS", 10), getenv("BS", 512), high=X_train.shape[0])
+  X_samp, Y_samp = X_train[samples], Y_train[samples]
+  print("*** got samples")
+
+  with Tensor.train():
+    # TODO: this shouldn't be a for loop. something like: (contract is still up in the air)
+    """
+    i = UOp.range(samples.shape[0])  # TODO: fix range function on UOp
+    losses = model(X_samp[i]).sparse_categorical_crossentropy(Y_samp[i]).backward().contract(i)
+    opt.schedule_steps(i)
+    """
+    losses = []
+    for i in range(samples.shape[0]):
+      opt.zero_grad()
+      losses.append(model(X_samp[i]).sparse_categorical_crossentropy(Y_samp[i]).backward())
+      opt.schedule_step()
+    # TODO: this stack currently breaks the "generator" aspect of losses. it probably shouldn't
+    #losses = Tensor.stack(*losses)
+  print("*** scheduled training")
+
+  # evaluate the model
+  with Tensor.test():
+    test_acc = ((model(X_test).argmax(axis=1) == Y_test).mean()*100)
+  print("*** scheduled eval")
+
+  # NOTE: there's no kernels run in the scheduling phase
+  assert GlobalCounters.kernel_count == 0, "kernels were run during scheduling!"
+
+  # only actually do anything at the end
+  if getenv("LOSS", 1):
+    for i in (t:=trange(len(losses))): t.set_description(f"loss: {losses[i].item():6.2f}")
+  print(f"test_accuracy: {test_acc.item():5.2f}%")
@@ -1,6 +1,6 @@
 # thanks to https://github.com/openai/whisper for a good chunk of MIT licensed code

-import sys, base64, multiprocessing, itertools
+import sys, base64, multiprocessing, itertools, collections
 from typing import Optional, Union, Literal, List

 from tinygrad import Tensor, TinyJit, Variable, nn
@@ -101,24 +101,17 @@ class TextDecoder:
    self.blocks = [ResidualAttentionBlock(n_text_state, n_text_head, is_decoder_block=True, max_self_attn_cache_len=self.max_self_attn_cache_len) for _ in range(n_text_layer)]
    self.ln = nn.LayerNorm(n_text_state)
    self.mask = Tensor.full((n_text_ctx, n_text_ctx), -np.inf).triu(1).realize()
-    self.blocks_start_tok = [TinyJit(block.__call__) for block in self.blocks]
-    self.blocks_after_start_tok = [TinyJit(block.__call__) for block in self.blocks]
-    self.start_output_tok = TinyJit(self.output_tok)
-    self.after_start_output_tok = TinyJit(self.output_tok)
+    self.getjitted = collections.defaultdict(lambda: TinyJit(self.forward))

-  # if layernorm supported symbolic shapes, we wouldn't need this hacky 'streaming' param (which should be called something more descriptive like 'x_is_start_toks_only')
-  def __call__(self, x: Tensor, pos: int, encoded_audio: Tensor, streaming=False):
+  def __call__(self, x: Tensor, pos: int, encoded_audio: Tensor):
+    pos = Variable("self_attn_cache_len", 1, self.max_self_attn_cache_len).bind(pos) if pos else 0
+    return self.getjitted[x.shape](x, pos, encoded_audio)
+
+  def forward(self, x:Tensor, pos:Union[Variable, Literal[0]], encoded_audio:Tensor):
    seqlen = x.shape[-1]
-    x = self.token_embedding(x) + self.positional_embedding[pos:pos+seqlen]
-    if pos == 0:
-      for block in (self.blocks if streaming else self.blocks_start_tok):
-        x = block(x, xa=encoded_audio, mask=self.mask, len=0)  # pass xa for cross attn kv caching
-      return self.output_tok(x) if streaming else self.start_output_tok(x)
-    else:
-      for block in self.blocks_after_start_tok:
-        len_v = Variable("self_attn_cache_len", 1, self.max_self_attn_cache_len).bind(pos)
-        x = block(x, mask=self.mask, len=len_v)
-      return self.after_start_output_tok(x)
+    x = self.token_embedding(x) + self.positional_embedding.shrink(((pos, pos+seqlen), None, None))
+    for block in self.blocks: x = block(x, xa=encoded_audio, mask=self.mask, len=pos)
+    return self.output_tok(x)

  def output_tok(self, x):
    return (self.ln(x) @ self.token_embedding.weight.T).realize()
@@ -169,7 +162,7 @@ def prep_audio(waveforms: List[np.ndarray], batch_size: int, truncate=False) ->
  mel_spec = librosa.filters.mel(sr=RATE, n_fft=N_FFT, n_mels=N_MELS) @ magnitudes

  log_spec = np.log10(np.clip(mel_spec, 1e-10, None))
-  log_spec = np.maximum(log_spec, log_spec.max() - 8.0)
+  log_spec = np.maximum(log_spec, log_spec.max((1,2), keepdims=True) - 8.0)
  log_spec = (log_spec + 4.0) / 4.0

  return log_spec
@@ -243,19 +236,26 @@ def load_file_waveform(filename):
 def transcribe_file(model, enc, filename):
  return transcribe_waveform(model, enc, [load_file_waveform(filename)])

-def transcribe_waveform(model, enc, waveforms, truncate=False):
+def transcribe_waveform(model: Whisper, enc, waveforms, truncate=False):
  """
  Expects an array of shape (N,S) where N is the number waveforms to transcribe in parallel and S is number of 16000Hz samples
  Returns the transcribed text if a single waveform is provided, or an array of transcriptions if multiple are provided
  """
-  N_audio = len(waveforms)
+
  log_spec = prep_audio(waveforms, model.batch_size, truncate)
+  nsample = model.decoder.max_tokens_to_sample

-  if log_spec.shape[-1] > FRAMES_PER_SEGMENT and N_audio > 1:
-    # we don't support multi-segment batching because the size of the prompt tokens would be different for each item in the batch
-    # if we really want this feature, we can consider padding or trimming prompt tokens of varying lengths to make them consistent
-    raise Exception("Multi-segment transcription not supported with batch audio input")
+  def inferloop(ctx: Union[np.ndarray, List[np.ndarray]], encoded_audio):
+    pos, next_tokens = 0, ctx
+    for i in range((nsample-len(start_tokens))*2):
+      next_tokens = model.decoder(Tensor(next_tokens), pos, encoded_audio)[:, -1].argmax(axis=-1).numpy().astype(np.int32).reshape(-1, 1)
+      next_tokens[ctx[:, -1] == eot] = eot
+      ctx = np.concatenate((ctx, next_tokens), axis=1)
+      pos = ctx.shape[-1] - 1
+      if (next_tokens == eot).all(): break
+    return ctx

+  def gettexttoks(line): return [tok for tok in line if tok < eot or tok > enc._special_tokens["<|notimestamps|>"]][-nsample+len(start_tokens):]
  start_tokens = [enc._special_tokens["<|startoftranscript|>"]]
  if model.is_multilingual:
    # TODO detect language
@@ -263,40 +263,24 @@ def transcribe_waveform(model, enc, waveforms, truncate=False):
    start_tokens.append(language_token)
    start_tokens.append(enc._special_tokens["<|transcribe|>"])
  start_tokens.append(enc._special_tokens["<|notimestamps|>"])
-  transcription_start_index = len(start_tokens)
+
  eot = enc._special_tokens["<|endoftext|>"]
-  transcription_tokens = [np.array([], dtype=np.int32)] * log_spec.shape[0]
+
+  ctx = np.tile(start_tokens, (model.batch_size,1))
+  transcriptions = [[] for _ in waveforms]

  for curr_frame in range(0, log_spec.shape[-1], FRAMES_PER_SEGMENT):
    encoded_audio = model.encoder.encode(Tensor(log_spec[:, :, curr_frame:curr_frame + FRAMES_PER_SEGMENT]))
-    pos = 0
-    curr_segment_tokens = np.tile(start_tokens, (log_spec.shape[0], 1))
-    if curr_frame > 0:
-      # pass the previously inferred tokens as 'prompt' - https://github.com/openai/whisper/discussions/117#discussioncomment-3727051
-      prompt = np.concatenate((
-        [enc._special_tokens["<|startofprev|>"]],
-        transcription_tokens[0][-model.decoder.max_tokens_to_sample+1:],
-        start_tokens))
-      curr_segment_tokens = np.tile(prompt, (log_spec.shape[0], 1))
-      transcription_start_index = len(curr_segment_tokens[0])

-    for i in range(model.decoder.max_tokens_to_sample):
-      out = model.decoder(Tensor(curr_segment_tokens if i == 0 else curr_segment_tokens[:, -1:]), pos, encoded_audio, streaming=curr_frame > 0)
-      next_tokens = out[:, -1].argmax(axis=-1).numpy().astype(np.int32)
-      next_tokens[curr_segment_tokens[:, -1] == eot] = eot
-      curr_segment_tokens = np.concatenate((curr_segment_tokens, next_tokens.reshape(-1, 1)), axis=1)
-      pos = curr_segment_tokens.shape[-1] - 1
-      if DEBUG >= 1: print(i, list(map(lambda tokens: enc.decode(tokens), curr_segment_tokens)))
-      if (curr_segment_tokens[:, -1] == eot).all():
-        break
+    if all(len(c) == len(ctx[0]) for c in ctx): ctx = inferloop(np.array(ctx), encoded_audio)
+    else: ctx = [inferloop((np.array([c]*model.batch_size)), encoded_audio)[i] for i,c in enumerate(ctx)]

-    for i, t in enumerate(curr_segment_tokens):
-      eot_index = np.where(t == eot)[0]
-      eot_index = None if len(eot_index) == 0 else eot_index[0]
-      transcription_tokens[i] = np.concatenate((transcription_tokens[i], t[transcription_start_index:eot_index]))
+    for i, (res, arr) in enumerate(zip(transcriptions, ctx)):
+      if curr_frame*HOP_LENGTH <= len(waveforms[i]):res.extend(arr[np.where(arr == start_tokens[-1])[0][0]+1:eoti[0] if len (eoti:=np.where(arr == eot)[0]) else None])
+    ctx = [[enc._special_tokens['<|startofprev|>']]+gettexttoks(cs)+start_tokens for cs in ctx]

-  transcriptions = list(map(lambda tokens: enc.decode(tokens).strip(), transcription_tokens))
-  return transcriptions[:N_audio] if N_audio > 1 else transcriptions[0]
+  transcriptions = list(map(lambda tokens: enc.decode(tokens).strip(), transcriptions))
+  return transcriptions if len(transcriptions) > 1 else transcriptions[0]

 CHUNK = 1600
 RECORD_SECONDS = 10
@@ -90,7 +90,7 @@ def uops_to_triton(function_name:str, uops:List[UOp]):
      else: kk(f"{ssa(u, 'val')} = {render_cast(f'tl.where({r[vin[2]]}, tl.load({r[vin[0]]}+{fill_dims_for_idx(r[vin[1]],dims)} , mask={render_valid(valid+[r[vin[2]]])}), 0.0)', dtype)}")
    elif uop == UOps.DEFINE_ACC: kk(f"{ssa(u, 'acc')} = {define_scalar(local_size, dtype, args).replace('//', '/')}")
    elif uop == UOps.CONST: r[u] = define_scalar([], dtype, args)
-    elif uop == UOps.PHI:
+    elif uop == UOps.ASSIGN:
      kk(f"{r[vin[0]]} = {r[vin[1]].replace('//', '/')}")
      r[u] = r[vin[0]]
    elif uop == UOps.STORE:
@@ -15,25 +15,13 @@ BASEDIR = Path(__file__).parent / "kits19" / "data"
 TRAIN_PREPROCESSED_DIR =  Path(__file__).parent / "kits19" / "preprocessed" / "train"
 VAL_PREPROCESSED_DIR =  Path(__file__).parent / "kits19" / "preprocessed" / "val"

-"""
-To download the dataset:
-```sh
-git clone https://github.com/neheller/kits19
-cd kits19
-pip3 install -r requirements.txt
-python3 -m starter_code.get_imaging
-cd ..
-mv kits19 extra/datasets
-```
-"""
-
@functools.lru_cache(None)
 def get_train_files():
  return sorted([x for x in BASEDIR.iterdir() if x.stem.startswith("case") and int(x.stem.split("_")[-1]) < 210 and x not in get_val_files()])

@functools.lru_cache(None)
 def get_val_files():
-  data = fetch("https://raw.githubusercontent.com/mlcommons/training/master/image_segmentation/pytorch/evaluation_cases.txt").read_text()
+  data = fetch("https://raw.githubusercontent.com/mlcommons/training/master/retired_benchmarks/unet3d/pytorch/evaluation_cases.txt").read_text()
  return sorted([x for x in BASEDIR.iterdir() if x.stem.split("_")[-1] in data.split("\n")])

 def load_pair(file_path):
@@ -123,7 +111,7 @@ def pad_input(volume, roi_shape, strides, padding_mode="constant", padding_val=-
  paddings = [bounds[2]//2, bounds[2]-bounds[2]//2, bounds[1]//2, bounds[1]-bounds[1]//2, bounds[0]//2, bounds[0]-bounds[0]//2, 0, 0, 0, 0]
  return F.pad(torch.from_numpy(volume), paddings, mode=padding_mode, value=padding_val).numpy(), paddings

-def sliding_window_inference(model, inputs, labels, roi_shape=(128, 128, 128), overlap=0.5):
+def sliding_window_inference(model, inputs, labels, roi_shape=(128, 128, 128), overlap=0.5, gpus=None):
  from tinygrad.engine.jit import TinyJit
  mdl_run = TinyJit(lambda x: model(x).realize())
  image_shape, dim = list(inputs.shape[2:]), len(inputs.shape[2:])
@@ -152,7 +140,7 @@ def sliding_window_inference(model, inputs, labels, roi_shape=(128, 128, 128), o
  for i in range(0, strides[0] * size[0], strides[0]):
    for j in range(0, strides[1] * size[1], strides[1]):
      for k in range(0, strides[2] * size[2], strides[2]):
-        out = mdl_run(Tensor(inputs[..., i:roi_shape[0]+i,j:roi_shape[1]+j, k:roi_shape[2]+k])).numpy()
+        out = mdl_run(Tensor(inputs[..., i:roi_shape[0]+i,j:roi_shape[1]+j, k:roi_shape[2]+k], device=gpus)).numpy()
        result[..., i:roi_shape[0]+i, j:roi_shape[1]+j, k:roi_shape[2]+k] += out * norm_patch
        norm_map[..., i:roi_shape[0]+i, j:roi_shape[1]+j, k:roi_shape[2]+k] += norm_patch
  result /= norm_map
@@ -2,14 +2,13 @@ import sys
 import json
 import numpy as np
 from PIL import Image
-import pathlib
+from pathlib import Path
 import boto3, botocore
-from tinygrad.helpers import fetch
-from tqdm import tqdm
+from tinygrad.helpers import fetch, tqdm, getenv
 import pandas as pd
 import concurrent.futures

-BASEDIR = pathlib.Path(__file__).parent / "open-images-v6-mlperf"
+BASEDIR = Path(__file__).parent / "open-images-v6-mlperf"
 BUCKET_NAME = "open-images-dataset"
 TRAIN_BBOX_ANNOTATIONS_URL = "https://storage.googleapis.com/openimages/v6/oidv6-train-annotations-bbox.csv"
 VALIDATION_BBOX_ANNOTATIONS_URL = "https://storage.googleapis.com/openimages/v5/validation-annotations-bbox.csv"
@@ -55,17 +54,12 @@ MLPERF_CLASSES = ['Airplane', 'Antelope', 'Apple', 'Backpack', 'Balloon', 'Banan
 ]


-def openimages(subset: str):
+def openimages(base_dir:Path, subset:str, ann_file:Path):
  valid_subsets = ['train', 'validation']
  if subset not in valid_subsets:
    raise ValueError(f"{subset=} must be one of {valid_subsets}")

-  ann_file = BASEDIR / f"{subset}/labels/openimages-mlperf.json"
-
-  if not ann_file.is_file():
-    fetch_openimages(ann_file, subset)
-
-  return ann_file
+  fetch_openimages(ann_file, base_dir, subset)

 # this slows down the conversion a lot!
 # maybe use https://raw.githubusercontent.com/scardine/image_size/master/get_image_size.py
@@ -112,10 +106,10 @@ def download_image(bucket, subset, image_id, data_dir):
  except botocore.exceptions.ClientError as exception:
    sys.exit(f"ERROR when downloading image `validation/{image_id}`: {str(exception)}")

-def fetch_openimages(output_fn, subset: str):
+def fetch_openimages(output_fn:str, base_dir:Path, subset:str):
  bucket = boto3.resource("s3", config=botocore.config.Config(signature_version=botocore.UNSIGNED)).Bucket(BUCKET_NAME)

-  annotations_dir, data_dir = BASEDIR / "annotations", BASEDIR / f"{subset}/data"
+  annotations_dir, data_dir = base_dir / "annotations", base_dir / f"{subset}/data"
  annotations_dir.mkdir(parents=True, exist_ok=True)
  data_dir.mkdir(parents=True, exist_ok=True)

@@ -143,8 +137,8 @@ def fetch_openimages(output_fn, subset: str):
  print("Converting annotations to COCO format...")
  export_to_coco(class_map, annotations, image_list, data_dir, output_fn, subset)

-def image_load(subset, fn):
-  img_folder = BASEDIR / f"{subset}/data"
+def image_load(base_dir, subset, fn):
+  img_folder = base_dir / f"{subset}/data"
  img = Image.open(img_folder / fn).convert('RGB')
  import torchvision.transforms.functional as F
  ret = F.resize(img, size=(800, 800))
@@ -164,18 +158,28 @@ def prepare_target(annotations, img_id, img_size):
  classes = classes[keep]
  return {"boxes": boxes, "labels": classes, "image_id": img_id, "image_size": img_size}

-def iterate(coco, bs=8):
+def iterate(coco, base_dir, bs=8):
  image_ids = sorted(coco.imgs.keys())
  for i in range(0, len(image_ids), bs):
    X, targets  = [], []
    for img_id in image_ids[i:i+bs]:
      img_dict = coco.loadImgs(img_id)[0]
-      x, original_size = image_load(img_dict['subset'], img_dict["file_name"])
+      x, original_size = image_load(base_dir, img_dict['subset'], img_dict["file_name"])
      X.append(x)
      annotations = coco.loadAnns(coco.getAnnIds(img_id))
      targets.append(prepare_target(annotations, img_id, original_size))
    yield np.array(X), targets

+def download_dataset(base_dir:Path, subset:str) -> Path:
+  if (ann_file:=base_dir / f"{subset}/labels/openimages-mlperf.json").is_file(): print(f"{subset} dataset is already available")
+  else:
+    print(f"Downloading {subset} dataset...")
+    openimages(base_dir, subset, ann_file)
+    print("Done")
+
+  return ann_file
+
+
 if __name__ == "__main__":
-  openimages("validation")
-  openimages("train")
+  download_dataset(base_dir:=getenv("BASE_DIR", BASEDIR), "train")
+  download_dataset(base_dir, "validation")
@@ -1,6 +1,7 @@
 import ctypes
 import os
 import pathlib
+import struct
 from hexdump import hexdump

 fxn = None
@@ -14,7 +15,8 @@ def disasm_raw(buf):
  fxn(buf, len(buf))

 def disasm(buf):
-  END = b"\x00\x00\x00\x00\x00\x00\x00\x03"
-  buf = buf[0x510:]  # this right?
-  buf = buf.split(END)[0] + END
-  disasm_raw(buf)
+  def _read_lib(off): return struct.unpack("I", buf[off:off+4])[0]
+
+  image_offset = _read_lib(0xc0)
+  image_size = _read_lib(0x100)
+  disasm_raw(buf[image_offset:image_offset+image_size])
@@ -0,0 +1,125 @@
+#!/usr/bin/env python3
+import os, ctypes, time, fcntl, mmap
+import llvmlite.binding as llvm
+from tinygrad.helpers import getenv, to_mv
+from tinygrad.runtime.support.elf import elf_loader
+from hexdump import hexdump
+from tinygrad.runtime.autogen import libc
+if getenv("IOCTL"): import run # noqa: F401 # pylint: disable=unused-import
+
+adsp = ctypes.CDLL(ctypes.util.find_library("adsprpc"))
+import adsprpc
+import ion
+import msm_ion
+ION_IOC_ALLOC = 0
+ION_IOC_MAP = 2
+ION_IOC_SHARE = 4
+ION_IOC_CUSTOM = 6
+ION_ADSP_HEAP_ID = 22
+ION_IOMMU_HEAP_ID = 25
+
+def ion_iowr(fd, nr, args):
+  ret = fcntl.ioctl(fd, (3 << 30) | (ctypes.sizeof(args) & 0x1FFF) << 16 | (ord(ion.ION_IOC_MAGIC) & 0xFF) << 8 | (nr & 0xFF), args)
+  if ret != 0: raise RuntimeError(f"ioctl returned {ret}")
+
+if __name__ == "__main__":
+  # TODO: mmap tensors to the DSP
+  # call the target function with the mmaped tensors
+  ion_fd = os.open("/dev/ion", os.O_RDWR | os.O_CLOEXEC)
+  arg3 = ion.struct_ion_allocation_data(len=0x1000, align=0x1000, heap_id_mask=1<<msm_ion.ION_SYSTEM_HEAP_ID, flags=ion.ION_FLAG_CACHED)
+  ion_iowr(ion_fd, ION_IOC_ALLOC, arg3)
+  print(arg3.handle)
+  arg2 = ion.struct_ion_fd_data(handle=arg3.handle)
+  ion_iowr(ion_fd, ION_IOC_SHARE, arg2)
+  print(arg2.fd)
+
+  res = libc.mmap(0, 0x1000, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED, arg2.fd, 0)
+  print("mmapped", hex(res))
+  to_mv(res, 0x10)[1] = 0xaa
+
+  from tinygrad.runtime.ops_clang import ClangCompiler
+  cc = ClangCompiler(args=["--target=hexagon", "-mcpu=hexagonv65", "-fuse-ld=lld", "-nostdlib"])
+
+  obj = cc.compile("""
+    typedef unsigned long long remote_handle64;
+    typedef struct { void *pv; unsigned int len; } remote_buf;
+    typedef struct { int fd; unsigned int offset; } remote_dma_handle;
+    typedef union { remote_buf buf; remote_handle64 h64; remote_dma_handle dma; } remote_arg;
+    void* HAP_mmap(void *addr, int len, int prot, int flags, int fd, long offset);
+    int HAP_munmap(void *addr, int len);
+    #define HAP_MEM_CACHE_WRITETHROUGH 0x40
+
+    int entry(unsigned long long handle, unsigned int sc, remote_arg* pra) {
+      if (sc>>24 == 1) {
+        //void *mmaped = *((void**)pra[0].buf.pv);
+        void *a = HAP_mmap(0, 0x1000, 3, 0, pra[1].dma.fd, 0);
+        ((char*)a)[0] = 0x55;
+        ((char*)a)[4] = 0x55;
+        ((char*)a)[8] = 0x99;
+        //((char*)a)[1] = 0x9b;
+        //char ret = ((char*)a)[1];
+        HAP_munmap(a, 0x1000);
+        return 0;
+
+        //return ((int)mmaped)&0xFFFF;
+        //return ((char*)mmaped)[1];
+        //return sizeof(void*);
+
+        //((char*)mmaped)[0] = 55;
+        //return ((int)mmaped)&0xFFFF;
+
+        //void addr = *((void**)pra[1])
+        //return sizeof(remote_buf);
+        //((char*)pra[1].h64)[0] = 55;
+
+        //return ((char*)mmaped)[1];
+        //((char*)mmaped)[0] = 55;
+        // NOTE: you have to return 0 for outbufs to work
+        //return ((int)pra[1].h64)&0xFFFF;
+      }
+      return 0;
+    }
+  """)
+  with open("/tmp/swag.so", "wb") as f: f.write(obj)
+
+  handle = ctypes.c_int64(-1)
+  adsp.remote_handle64_open(ctypes.create_string_buffer(b"file:////tmp/swag.so?entry&_modver=1.0&_dom=cdsp"), ctypes.byref(handle))
+  print("HANDLE", handle.value)
+  #print(adsp.remote_handle64_invoke(handle, 0, None))
+
+  #rem = adsp.remote_register_buf(res, 0x1000, arg2.fd, 4)
+  #rem = adsp.remote_register_dma_handle(arg2.fd, 0x1000)
+  #print("remote_register_buf_attr", rem)
+
+  #out = ctypes.c_uint64(0)
+  #ret = adsp.remote_mmap(arg2.fd, 0, 0, 0x1000, ctypes.byref(out))
+  #print(ret)
+  #print("mapped at", hex(out.value))
+
+  #arg_2 = ctypes.c_int64(out.value)
+  arg_2 = ctypes.c_int64(arg2.fd)
+  pra = (adsprpc.union_remote_arg64 * 3)()
+  pra[0].buf.pv = ctypes.addressof(arg_2)
+  pra[0].buf.len = 8
+  pra[1].dma.fd = arg2.fd
+  pra[1].dma.len = 0x1000
+  print("invoke")
+  ret = adsp.remote_handle64_invoke(handle, (1<<24) | (1<<16) | (1 << 4), pra)
+  print("return value", ret, hex(ret))
+  #print(hex(arg_2.value), arg_2.value)
+  #time.sleep(0.1)
+
+  # flush the cache
+  """
+  flush_data = msm_ion.struct_ion_flush_data(handle=arg3.handle, vaddr=res, offset=0, length=0x1000)
+  # ION_IOC_CLEAN_INV_CACHES
+  cd = ion.struct_ion_custom_data(
+    cmd=(3 << 30) | (ctypes.sizeof(flush_data) & 0x1FFF) << 16 | (ord(msm_ion.ION_IOC_MSM_MAGIC) & 0xFF) << 8 | (2 & 0xFF),
+    arg=ctypes.addressof(flush_data))
+  ret = ion_iowr(ion_fd, ION_IOC_CUSTOM, cd)
+
+  res2 = libc.mmap(0, 0x1000, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED, arg2.fd, 0)
+  """
+  hexdump(to_mv(res, 0x10))
+  os._exit(0)
+
@@ -0,0 +1,3 @@
+#!/bin/bash
+clang2py adsprpc_shared.h -k cdefstum -o adsprpc.py
+
@@ -0,0 +1,328 @@
+/*
+ * Copyright (c) 2005-2007, 2012-2013, 2019-2020 Qualcomm Technologies, Inc.
+ * All Rights Reserved.
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its contributors
+ * may be used to endorse or promote products derived from this software without
+ * specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+ * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+ * POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef AEESTDERR_H
+#define AEESTDERR_H
+//
+// Basic Error Codes
+//
+//
+#if defined(__hexagon__)
+	#define AEE_EOFFSET               0x80000400
+#else
+	#define AEE_EOFFSET               0x00000000
+#endif
+/** @defgroup stdbasicerror Basic error codes
+ *  @{
+ */
+#define AEE_SUCCESS                   0                      ///< No error
+#define AEE_EUNKNOWN                  -1                     ///< Unknown error (should not use this)
+
+#define AEE_EFAILED                   (AEE_EOFFSET + 0x001)  ///< General failure
+#define AEE_ENOMEMORY                 (AEE_EOFFSET + 0x002)  ///< Memory allocation failed because of insufficient RAM
+#define AEE_ECLASSNOTSUPPORT          (AEE_EOFFSET + 0x003)  ///< Specified class unsupported
+#define AEE_EVERSIONNOTSUPPORT        (AEE_EOFFSET + 0x004)  ///< Version not supported
+#define AEE_EALREADYLOADED            (AEE_EOFFSET + 0x005)  ///< Object already loaded
+#define AEE_EUNABLETOLOAD             (AEE_EOFFSET + 0x006)  ///< Unable to load object/applet
+#define AEE_EUNABLETOUNLOAD           (AEE_EOFFSET + 0x007)  ///< Unable to unload
+                                                                    ///< object/applet
+#define AEE_EALARMPENDING             (AEE_EOFFSET + 0x008)  ///< Alarm is pending
+#define AEE_EINVALIDTIME              (AEE_EOFFSET + 0x009)  ///< Invalid time
+#define AEE_EBADCLASS                 (AEE_EOFFSET + 0x00A)  ///< NULL class object
+#define AEE_EBADMETRIC                (AEE_EOFFSET + 0x00B)  ///< Invalid metric specified
+#define AEE_EEXPIRED                  (AEE_EOFFSET + 0x00C)  ///< App/Component Expired
+#define AEE_EBADSTATE                 (AEE_EOFFSET + 0x00D)  ///< Process or thread is not in expected state
+#define AEE_EBADPARM                  (AEE_EOFFSET + 0x00E)  ///< Invalid parameter
+#define AEE_ESCHEMENOTSUPPORTED       (AEE_EOFFSET + 0x00F)  ///< Invalid URL scheme
+#define AEE_EBADITEM                  (AEE_EOFFSET + 0x010)  ///< Value out of range
+#define AEE_EINVALIDFORMAT            (AEE_EOFFSET + 0x011)  ///< Invalid format
+#define AEE_EINCOMPLETEITEM           (AEE_EOFFSET + 0x012)  ///< Incomplete item, like length of a string is less that expected
+#define AEE_ENOPERSISTMEMORY          (AEE_EOFFSET + 0x013)  ///< Insufficient flash
+#define AEE_EUNSUPPORTED              (AEE_EOFFSET + 0x014)  ///< API not implemented
+#define AEE_EPRIVLEVEL                (AEE_EOFFSET + 0x015)  ///< Privileges are insufficient
+                                                                    ///< for this operation
+#define AEE_ERESOURCENOTFOUND         (AEE_EOFFSET + 0x016)  ///< Unable to find specified
+                                                                    ///< resource
+#define AEE_EREENTERED                (AEE_EOFFSET + 0x017)  ///< Non re-entrant API
+                                                                    ///< re-entered
+#define AEE_EBADTASK                  (AEE_EOFFSET + 0x018)  ///< API called in wrong task
+                                                                    ///< context
+#define AEE_EALLOCATED                (AEE_EOFFSET + 0x019)  ///< App/Module left memory
+                                                                    ///< allocated when released.
+#define AEE_EALREADY                  (AEE_EOFFSET + 0x01A)  ///< Operation is already in
+                                                                    ///< progress
+#define AEE_EADSAUTHBAD               (AEE_EOFFSET + 0x01B)  ///< ADS mutual authorization
+                                                                    ///< failed
+#define AEE_ENEEDSERVICEPROG          (AEE_EOFFSET + 0x01C)  ///< Need service programming
+#define AEE_EMEMPTR                   (AEE_EOFFSET + 0x01D)  ///< bad memory pointer, expected to be NULL
+#define AEE_EHEAP                     (AEE_EOFFSET + 0x01E)  ///< An internal heap error was detected
+#define AEE_EIDLE                     (AEE_EOFFSET + 0x01F)  ///< Context (system, interface,
+                                                                    ///< etc.) is idle
+#define AEE_EITEMBUSY                 (AEE_EOFFSET + 0x020)  ///< Context (system, interface,
+                                                                    ///< etc.) is busy
+#define AEE_EBADSID                   (AEE_EOFFSET + 0x021)  ///< Invalid subscriber ID
+#define AEE_ENOTYPE                   (AEE_EOFFSET + 0x022)  ///< No type detected/found
+#define AEE_ENEEDMORE                 (AEE_EOFFSET + 0x023)  ///< Need more data/info
+#define AEE_EADSCAPS                  (AEE_EOFFSET + 0x024)  ///< ADS Capabilities do not
+                                                                    ///< match those required for phone
+#define AEE_EBADSHUTDOWN              (AEE_EOFFSET + 0x025)  ///< App failed to close properly
+#define AEE_EBUFFERTOOSMALL           (AEE_EOFFSET + 0x026)  ///< Destination buffer given is
+                                                                    ///< too small
+                                                                    ///< or service exists or is
+                                                                    ///< valid
+#define AEE_EACKPENDING               (AEE_EOFFSET + 0x028)  ///< ACK pending on application
+#define AEE_ENOTOWNER                 (AEE_EOFFSET + 0x029)  ///< Not an owner authorized to
+                                                                    ///< perform the operation
+#define AEE_EINVALIDITEM              (AEE_EOFFSET + 0x02A)  ///< Current item is invalid, it can be a switch case or a pointer to memory
+#define AEE_ENOTALLOWED               (AEE_EOFFSET + 0x02B)  ///< Not allowed to perform the
+                                                                    ///< operation
+#define AEE_EBADHANDLE                (AEE_EOFFSET + 0x02C)  ///< Invalid/Wrong handle
+#define AEE_EINVHANDLE                (AEE_EOFFSET + 0x02C)  ///< Invalid handle - adding here as its defined in vendor AEEStdErr.h - needed to check valid handle in stub.c
+#define AEE_EOUTOFHANDLES             (AEE_EOFFSET + 0x02D)  ///< Out of handles (Handle list is already full)
+//Hole here
+#define AEE_ENOMORE                   (AEE_EOFFSET + 0x02F)  ///< No more items available --
+                                                                    ///< reached end
+#define AEE_ECPUEXCEPTION             (AEE_EOFFSET + 0x030)  ///< A CPU exception occurred
+#define AEE_EREADONLY                 (AEE_EOFFSET + 0x031)  ///< Cannot change read-only
+                                                                    ///< object or parameter ( Parameter is in protected mode)
+#define AEE_ERPC                      (AEE_EOFFSET + 0x200)  ///< Error due to fastrpc implementation
+#define AEE_EFILE                     (AEE_EOFFSET + 0x201)  ///<File handling related error
+//NOTE: Used in both HLOS and DSP.
+#define AEE_ENOSUCH                   (39)                   ///< No such name, port, socket
+#define AEE_EINTERRUPTED              (46)                   ///< Waitable call is interrupted,
+                                                                   ///< the user should return to the HLOS and retry the call
+#define AEE_ECONNRESET                (104)                  ///< Connection reset by peer
+#define AEE_EWOULDBLOCK               (516)                  ///< Operation would block if not
+                                                                    ///< non-blocking; wait and try
+                                                                    ///< again
+/**
+ * @}
+ */
+
+/** @defgroup sigverifyerror Sigverify error codes
+ *  @{
+ */
+
+#define AEE_EINVALIDMSG	              (AEE_EOFFSET + 0x032)     ///<  Invalid SMD message from APPS
+#define AEE_EINVALIDTHREAD            (AEE_EOFFSET + 0x033)     ///<  Invalid thread
+#define AEE_EINVALIDPROCESS           (AEE_EOFFSET + 0x034)     ///<  Invalid Process
+#define AEE_EINVALIDFILENAME          (AEE_EOFFSET + 0x035)     ///<  Invalid filename
+#define AEE_EINVALIDDIGESTSIZE        (AEE_EOFFSET + 0x036)     ///<  Invalid digest size
+#define AEE_EINVALIDSEGS              (AEE_EOFFSET + 0x037)     ///<  Invalid segments
+#define AEE_EINVALIDSIGNATURE         (AEE_EOFFSET + 0x038)     ///<  Invalid signature
+#define AEE_EINVALIDDOMAIN            (AEE_EOFFSET + 0x039)     ///<  Invalid DSP domain
+#define AEE_EINVALIDFD                (AEE_EOFFSET + 0x03A)     ///<  Invalid file descriptor
+#define AEE_EINVALIDDEVICE            (AEE_EOFFSET + 0x03B)     ///<  Invalid Device or Device node open failed for the domain
+#define AEE_EINVALIDMODE              (AEE_EOFFSET + 0x03C)     ///<  Invalid Mode
+#define AEE_EINVALIDPROCNAME          (AEE_EOFFSET + 0x03D)     ///<  Invalid Process name
+#define AEE_ENOSUCHMOD                (AEE_EOFFSET + 0x03E)     ///<  No such module
+#define AEE_ENOSUCHINSTANCE           (AEE_EOFFSET + 0x03F)     ///<  No instance in the list lookup
+#define AEE_ENOSUCHTHREAD             (AEE_EOFFSET + 0x040)     ///<  No such thread
+#define AEE_ENOSUCHPROCESS            (AEE_EOFFSET + 0x041)     ///<  No such process
+#define AEE_ENOSUCHSYMBOL             (AEE_EOFFSET + 0x042)     ///<  No such symbol( dlsym for the symbol failed)
+#define AEE_ENOSUCHDEVICE             (AEE_EOFFSET + 0x043)     ///<  No such device
+#define AEE_ENOSUCHPROP               (AEE_EOFFSET + 0x044)     ///<  No such dal property
+#define AEE_ENOSUCHFILE               (AEE_EOFFSET + 0x045)     ///<  No such file found
+#define AEE_ENOSUCHHANDLE             (AEE_EOFFSET + 0x046)     ///<  No such handle
+#define AEE_ENOSUCHSTREAM             (AEE_EOFFSET + 0x047)     ///<  No such stream
+#define AEE_ENOSUCHMAP                (AEE_EOFFSET + 0x048)     ///<  No mapping exists for this address on DSP
+#define AEE_ENOSUCHREGISTER           (AEE_EOFFSET + 0x049)     ///<  No such register
+#define AEE_ENOSUCHCLIENT             (AEE_EOFFSET + 0x04A)     ///<  No such QDI client
+#define AEE_EBADDOMAIN                (AEE_EOFFSET + 0x04B)     ///<  Bad domain (not initialized)
+#define AEE_EBADOFFSET                (AEE_EOFFSET + 0x04C)     ///<  Bad buffer/page/heap offset
+#define AEE_EBADSIZE                  (AEE_EOFFSET + 0x04D)     ///<  Bad buffer/page/heap size
+#define AEE_EBADPERMS                 (AEE_EOFFSET + 0x04E)     ///<  Bad FILE/MAP/MEM permissions
+#define AEE_EBADFD                    (AEE_EOFFSET + 0x04F)     ///<  Bad file descriptor
+#define AEE_EBADPID                   (AEE_EOFFSET + 0x050)     ///<  Bad PID from HLOS
+#define AEE_EBADTID                   (AEE_EOFFSET + 0x051)     ///<  Bad TID
+#define AEE_EBADELF                   (AEE_EOFFSET + 0x052)     ///<  Bad elf file
+#define AEE_EBADASID                  (AEE_EOFFSET + 0x053)     ///<  Bad asid
+#define AEE_EBADCONTEXT               (AEE_EOFFSET + 0x054)     ///<  Bad context
+#define AEE_EBADMEMALIGN              (AEE_EOFFSET + 0x055)     ///<  Bad memory alignment
+#define AEE_EIOCTL                    (AEE_EOFFSET + 0x056)     ///<  ioctl call failed
+#define AEE_EFOPEN                    (AEE_EOFFSET + 0x057)     ///<  file open error or device node open failed for DSP domain
+#define AEE_EFGETS                    (AEE_EOFFSET + 0x058)     ///<  file get string error
+#define AEE_EFFLUSH                   (AEE_EOFFSET + 0x059)     ///<  file flush error
+#define AEE_EFCLOSE                   (AEE_EOFFSET + 0x05A)     ///<  file close error
+#define AEE_EEOF                      (AEE_EOFFSET + 0x05B)     ///<  File EOF reached
+#define AEE_EFREAD                    (AEE_EOFFSET + 0x05C)     ///<  file read failed
+#define AEE_EFWRITE                   (AEE_EOFFSET + 0x05D)     ///<  file write failed
+#define AEE_EFGETPOS                  (AEE_EOFFSET + 0x05E)     ///<  file get position failed
+#define AEE_EFSETPOS                  (AEE_EOFFSET + 0x05F)     ///<  file set position failed
+#define AEE_EFTELL                    (AEE_EOFFSET + 0x060)     ///<  file tell position failed
+#define AEE_EFSEEK                    (AEE_EOFFSET + 0x061)     ///<  file seek failed
+#define AEE_EFLEN                     (AEE_EOFFSET + 0x062)     ///<  file len greater than expected
+#define AEE_EGETENV                   (AEE_EOFFSET + 0x063)     ///<  apps_std get enviroment failed
+#define AEE_ESETENV                   (AEE_EOFFSET + 0x064)     ///<  apps_std set enviroment failed
+#define AEE_EMMAP                     (AEE_EOFFSET + 0x065)     ///<  mmap failed
+#define AEE_EIONMAP                   (AEE_EOFFSET + 0x066)     ///<  ion map failed
+#define AEE_EIONALLOC                 (AEE_EOFFSET + 0x067)     ///<  ion alloc failed
+#define AEE_ENORPCMEMORY              (AEE_EOFFSET + 0x068)     ///<  ION memory allocation failed
+#define AEE_ENOROOTOFTRUST            (AEE_EOFFSET + 0x069)     ///<  No root of trust for sigverify
+#define AEE_ENOTLOCKED                (AEE_EOFFSET + 0x06A)     ///<  Unlock failed, not locked before
+#define AEE_ENOTINITIALIZED           (AEE_EOFFSET + 0x06B)     ///<  Not initialized
+#define AEE_EUNSUPPORTEDAPI           (AEE_EOFFSET + 0x06C)     ///<  unsupported API/request ID
+#define AEE_EUNPACK                   (AEE_EOFFSET + 0x06D)     ///<  unpacking command failed
+#define AEE_EPOLL                     (AEE_EOFFSET + 0x06E)     ///<  error while polling for event
+#define AEE_EEVENTREAD                (AEE_EOFFSET + 0x06F)     ///<  event read failed
+#define AEE_EMAXBUFS                  (AEE_EOFFSET + 0x070)     ///<  Maximum buffers
+#define AEE_EINVARGS                  (AEE_EOFFSET + 0x071)     ///<  Invalid Arguments
+#define AEE_ECONNREFUSED              (AEE_EOFFSET + 0x072)     ///<  Connection refused to DSP
+#define AEE_EUNSIGNEDMOD              (AEE_EOFFSET + 0x081)     ///<  test-sig not found, Unsigned shared object
+#define AEE_EINVALIDHASH              (AEE_EOFFSET + 0x082)     ///<  test-sig not found, Invalid hash object
+#define AEE_EBADVA                    (AEE_EOFFSET + 0x083)     ///<  Bad VA address
+#define AEE_ENOSUCHJOB                (AEE_EOFFSET + 0x084)     ///<  No such job
+#define AEE_ENOSUCHGROUP              (AEE_EOFFSET + 0x084)     ///<  No such static pd group
+#define AEE_EBADMAPREFCNT             (AEE_EOFFSET + 0x085)     ///<  Bad map reference count
+#define AEE_EBADPAGECNT               (AEE_EOFFSET + 0x086)     ///<  Bad page count
+#define AEE_EMAPALREADYPRESENT        (AEE_EOFFSET + 0x087)     ///<  Map already present
+#define AEE_ENOFREESECTION            (AEE_EOFFSET + 0x088)     ///<  No more free sections available
+#define AEE_U2GCLIENT_OPEN            (AEE_EOFFSET + 0x089)     ///<  u2g client open failed
+
+/**
+ * @}
+ */
+
+/** @defgroup smderror SMD error codes
+ *  @{
+ */
+
+
+#if defined(__hexagon__)
+	#define AEE_EGLINK_OFFSET         (AEE_EOFFSET + 0x100)     ///<  SMD errors offset
+	#define AEE_EGLINKBADPACKET       (AEE_EOFFSET + 0x101)     ///<  SMD invalid packet size
+	#define AEE_EGLINKALREADYOPEN     (AEE_EOFFSET + 0x102)     ///<  SMD port is already open
+	#define AEE_EGLINKOPENFAILED      (AEE_EOFFSET + 0x103)     ///<  SMD port open failed
+	#define AEE_EGLINKWRITE           (AEE_EOFFSET + 0x104)     ///<  SMD port write failed
+	#define AEE_EGLINKREGISTER        (AEE_EOFFSET + 0x105)     ///<  SMD port register callback failed
+#else
+	#define AEE_ESMD_OFFSET           (AEE_EOFFSET + 0x100)     ///<  SMD errors offset
+	#define AEE_ESMDBADPACKET         (AEE_EOFFSET + 0x101)     ///<  SMD invalid packet size
+	#define AEE_ESMDALREADYOPEN       (AEE_EOFFSET + 0x102)     ///<  SMD port is already open
+	#define AEE_ESMDOPENFAILED        (AEE_EOFFSET + 0x103)     ///<  SMD port open failed
+#endif
+ /**
+ * @}
+ */
+
+ /** @defgroup dalerror DAL error codes
+ *  @{
+ */
+
+
+#define AEE_EDAL_OFFSET               (AEE_EOFFSET + 0x120)     ///<  Dal error offset
+#define AEE_EDALDEVATTACH             (AEE_EOFFSET + 0x121)     ///<  DAL attach error
+#define AEE_EDALINTREGISTER           (AEE_EOFFSET + 0x122)     ///<  DAL interrupt register error
+#define AEE_EDALINTUNREGISTER         (AEE_EOFFSET + 0x123)     ///<  Dal interrupt unregister error
+#define AEE_EDALGETPROP               (AEE_EOFFSET + 0x124)     ///<  Dal get property
+#define AEE_EDALGETVAL                (AEE_EOFFSET + 0x125)     ///<  Dal get property value
+#define AEE_EDCVSREQUEST              (AEE_EOFFSET + 0x126)     ///<  Dal get property value
+
+ /**
+ * @}
+ */
+
+ /** @defgroup qurterror QURT error codes
+ *  @{
+ */
+
+#define AEE_EQURT_OFFSET              (AEE_EOFFSET + 0x140)     ///<  QURT error offset
+#define AEE_EQURTREGIONCREATE         (AEE_EOFFSET + 0x141)     ///<  QURT region create failed
+#define AEE_EQURTCACHECLEAN	          (AEE_EOFFSET + 0x142)     ///<  QURT cache clean failed
+#define AEE_EQURTREGIONGETATTR        (AEE_EOFFSET + 0x143)     ///<  QURT region get attribute failed
+#define AEE_EQURTBADREGIONPERMS       (AEE_EOFFSET + 0x144)     ///<  QURT bad permissions for region
+#define AEE_EQURTMEMPOOLADD	          (AEE_EOFFSET + 0x145)     ///<  QURT Add to memory pool failed
+#define AEE_EQURTREGISTERDEV          (AEE_EOFFSET + 0x146)     ///<  QURT register device failed
+#define AEE_EQURTMEMPOOLCREATE        (AEE_EOFFSET + 0x147)     ///<  QURT create memory pool failed
+#define AEE_EQURTGETVA                (AEE_EOFFSET + 0x148)     ///<  QURT get VA failed
+#define AEE_EQURTREGIONDELETE         (AEE_EOFFSET + 0x149)     ///<  QURT region delete failed
+#define AEE_EQURTMEMPOOLATTACH        (AEE_EOFFSET + 0x14A)     ///<  QURT memory pool attach failed
+#define AEE_EQURTTHREADCREATE         (AEE_EOFFSET + 0x14B)     ///<  QURT thread create failed
+#define AEE_EQURTCOPYTOUSER           (AEE_EOFFSET + 0x14C)     ///<  QURT copy to user memory failed
+#define AEE_EQURTMEMMAPCREATE         (AEE_EOFFSET + 0x14D)     ///<  QURT map create failed
+#define AEE_EQURTINVHANDLE            (AEE_EOFFSET + 0x14E)     ///<  QURT Invalid client handle
+#define AEE_EQURTBADASID              (AEE_EOFFSET + 0x14F)     ///<  QURT Bad ASIC from QURT
+#define AEE_EQURTOPENFAILED           (AEE_EOFFSET + 0x150)     ///<  QURT QDI open failed
+#define AEE_EQURTCOPYFROMUSER         (AEE_EOFFSET + 0x151)     ///<  QURT Copy from user failed
+#define AEE_EQURTLINELOCK             (AEE_EOFFSET + 0x152)     ///<  QURT Line lock failed
+#define AEE_EQURTQDIDEFMETHOD         (AEE_EOFFSET + 0x153)     ///<  QURT QDI default method failed
+#define AEE_EQURTCREATEHANDLE         (AEE_EOFFSET + 0x154)     ///<  QURT create handle from obj failed
+#define AEE_EQURTWRITABLEMEM          (AEE_EOFFSET + 0x155)     ///<  QURT CPZ migration writable mem
+#define AEE_EQURTTHREADCREATEDEF      (AEE_EOFFSET + 0x156)     ///<  QURT thread create def
+#define AEE_EQURTLOOKUPVA             (AEE_EOFFSET + 0x157)     ///<  QURT lookup VA
+#define AEE_EQURTLOOKUPPA             (AEE_EOFFSET + 0x158)     ///<  QURT lookup PA
+#define AEE_EQURTMIGRATESECURE        (AEE_EOFFSET + 0x159)     ///<  QURT CPZ migration failure
+#define AEE_EQURTQDIOPEN              (AEE_EOFFSET + 0X160)     ///<  QURT QDI open failure
+#define AEE_EQURTMAPREMOVE            (AEE_EOFFSET + 0X161)     ///<  QURT map remove failure
+#define AEE_EQURTQDICLOSE             (AEE_EOFFSET + 0X162)     ///<  QURT QDI close failed
+#define AEE_EQURTWAIT                 (AEE_EOFFSET + 0X163)     ///<  QURT Futex wait failed
+
+ /**
+ * @}
+ */
+
+  /** @defgroup mmpmerr MMPM error codes
+ *  @{
+ */
+
+#define AEE_EMMPM_OFFSET              (AEE_EOFFSET + 0x170)     ///<  MMPM errors offset
+#define AEE_EMMPMREQUEST              (AEE_EOFFSET + 0x171)     ///<  MMPM Power request to failed
+#define AEE_EMMPMRELEASE              (AEE_EOFFSET + 0x172)     ///<  MMPM Release request failed
+#define AEE_EMMPMSETPARAM             (AEE_EOFFSET + 0x173)     ///<  MMPM set param request failed
+#define AEE_EMMPMREGISTER             (AEE_EOFFSET + 0x174)     ///<  MMPM Register request failed
+#define AEE_EMMPMGETINFO              (AEE_EOFFSET + 0x175)     ///<  MMPM Get info failed
+#define AEE_EMAX_MMPM_CLIENTS         (AEE_EOFFSET + 0x176)     ///<  MMPM Reached maximum clients per PD(HAP_MAX_CLIENTS)
+#define AEE_EDCVSREGISTER             (AEE_EOFFSET + 0x177)     ///<  ADSP DCVS client registration failed
+#define AEE_PDRREGFAIL                (AEE_EOFFSET + 0x178)     ///<  Error Callback Services Registration failed for PD
+
+ /**
+ * @}
+ */
+
+#define AEE_DEFAULT_PROCESS           (AEE_EOFFSET + 0x180)     ///<  Default process in Guest OS is not present
+#define AEE_ENULLCONTEXT              (AEE_EOFFSET + 0x181)     ///<  User NULL context vote
+#define AEE_EINVALIDJOB               (AEE_EOFFSET + 0x182)     ///<  AsyncRPC Invalid job
+#define AEE_EBUSY                     (AEE_EOFFSET + 0x183)     ///<  AsyncRPC Pending job
+
+ /** @defgroup heaperror Heap error codes
+ *  @{
+ */
+
+#define E_APPS_BUSY_RETRY_LATER       (AEE_EOFFSET + 0x190)     ///<  Retry because the apps is busy
+#define E_HLOS_CAP_REACHED            (AEE_EOFFSET + 0x191)     ///<  cannot allocate any more hlos mem
+#define E_DPOOL_CAP_REACHED           (AEE_EOFFSET + 0x192)     ///<  cannot allocate any more physpool mem
+#define E_NO_MORE_FREE_SECTIONS       (AEE_EOFFSET + 0x193)     ///<  No more free sections available to grow heap
+
+ /**
+ * @}
+ */
+
+#endif /* #ifndef AEESTDERR_H */
+
@@ -0,0 +1,685 @@
+/*==============================================================================
+@file
+    HAP_power.h
+
+@brief
+    Header file of DSP power APIs.
+
+Copyright (c) 2015,2019 Qualcomm Technologies, Inc.
+All rights reserved. Qualcomm Proprietary and Confidential.
+==============================================================================*/
+
+#ifndef _HAP_POWER_H
+#define _HAP_POWER_H
+
+#include "AEEStdErr.h"
+//#include <string.h>
+//#include <stdlib.h>
+#define boolean char
+#define FALSE 0
+#define TRUE 1
+#define uint64 unsigned long long
+#define uint32 unsigned int
+#define NULL 0
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+//Add a weak reference so shared objects do not throw link error
+#pragma weak HAP_power_destroy_client
+
+/**
+* Possible error codes returned
+*/
+typedef enum {
+	HAP_POWER_ERR_UNKNOWN           = -1,
+	HAP_POWER_ERR_INVALID_PARAM     = -2,
+	HAP_POWER_ERR_UNSUPPORTED_API   = -3
+} HAP_power_error_codes;
+
+/** Payload for HAP_power_set_mips_bw */
+typedef struct {
+	boolean set_mips;						/**< Set to TRUE to request MIPS */
+	unsigned int mipsPerThread;				/**< mips requested per thread, to establish a minimal clock frequency per HW thread */
+	unsigned int mipsTotal;					/**< Total mips requested, to establish total number of MIPS required across all HW threads */
+	boolean set_bus_bw;						/**< Set to TRUE to request bus_bw */
+	uint64 bwBytePerSec;					/**< Max bus BW requested (bytes per second) */
+	unsigned short busbwUsagePercentage;	/**< Percentage of time during which bwBytesPerSec BW is required from the bus (0..100) */
+	boolean set_latency;					/**< Set to TRUE to set latency */
+	int latency;							/**< maximum hardware wakeup latency in microseconds.  The higher the value,
+											*	the deeper state of sleep that can be entered but the longer it may take
+											*	to awaken. Only values > 0 are supported (1 microsecond is the smallest valid value) */
+} HAP_power_mips_bw_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+ /** Clock frequency match type*/
+typedef enum {
+	HAP_FREQ_AT_LEAST,				/**< Matches at least the specified frequency. */
+	HAP_FREQ_AT_MOST,				/**< Matches at most the specified frequency. */
+	HAP_FREQ_CLOSEST,				/**< Closest match to the specified frequency. */
+	HAP_FREQ_EXACT,					/**< Exact match with the specified frequency. */
+	HAP_FREQ_MAX_COUNT				/**< Maximum count. */
+} HAP_freq_match_type;
+/**
+ * @} // HAP_power_enums
+ */
+
+/** Configuration for bus bandwidth */
+typedef struct {
+	boolean set_bus_bw;						/**< Set to TRUE to request bus_bw */
+	uint64 bwBytePerSec;					/**< Max bus BW requested (bytes per second) */
+	unsigned short busbwUsagePercentage;	/**< Percentage of time during which bwBytesPerSec BW is required from the bus (0..100) */
+} HAP_power_bus_bw;
+
+/**
+* @brief Payload for vapps power request
+* vapps core is used for Video post processing
+*/
+typedef struct {
+	boolean set_clk;						/**< Set to TRUE to request clock frequency */
+	unsigned int clkFreqHz;					/**< Clock frequency in Hz */
+	HAP_freq_match_type freqMatch;			/**< Clock frequency match */
+	HAP_power_bus_bw dma_ext;				/**< DMA external bus bandwidth */
+	HAP_power_bus_bw hcp_ext;				/**< HCP external bus bandwidth */
+	HAP_power_bus_bw dma_int;				/**< DMA internal bus bandwidth */
+	HAP_power_bus_bw hcp_int;				/**< HCP internal bus bandwidth */
+} HAP_power_vapss_payload;
+
+/**
+* @brief Payload for vapps_v2 power request
+* Supported in targets which have split VAPPS core(DMA and HCP) form Hana onwards
+*/
+typedef struct {
+	boolean set_dma_clk;					/**< Set to TRUE to reqeust DMA clock frequency */
+	boolean set_hcp_clk;					/**< Set to TRUE to reqeust HCP clock frequency */
+	unsigned int dmaClkFreqHz;				/**< DMA Clock frequency in Hz */
+	unsigned int hcpClkFreqHz;				/**< HCP Clock frequency in Hz */
+	HAP_freq_match_type freqMatch;			/**< Clock frequency match type */
+	HAP_power_bus_bw dma_ext;				/**< DMA external bus bandwidth */
+	HAP_power_bus_bw hcp_ext;				/**< HCP external bus bandwidth */
+	HAP_power_bus_bw dma_int;				/**< DMA internal bus bandwidth */
+	HAP_power_bus_bw hcp_int;				/**< HCP internal bus bandwidth */
+} HAP_power_vapss_payload_v2;
+
+/** Payload for HAP_power_set_HVX */
+typedef struct {
+	boolean power_up;						/**< Set to TRUE to turn on HVX, and FALSE to turn off. */
+} HAP_power_hvx_payload;
+
+/**
+* Payload for HAP_power_set_HMX
+* Supported from Lahaina onwards*/
+typedef struct {
+	boolean power_up;						/**< Set to TRUE to turn on HMX, and FALSE to turn off. */
+} HAP_power_hmx_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+/** Payload for HAP power client classes */
+typedef enum {
+	HAP_POWER_UNKNOWN_CLIENT_CLASS			= 0x00,		/**< Unknown client class */
+	HAP_POWER_AUDIO_CLIENT_CLASS			= 0x01,		/**< Audio client class */
+	HAP_POWER_VOICE_CLIENT_CLASS			= 0x02,		/**< Voice client class */
+	HAP_POWER_COMPUTE_CLIENT_CLASS			= 0x04,		/**< Compute client class */
+	HAP_POWER_STREAMING_1HVX_CLIENT_CLASS	= 0x08,		/**< Camera streaming with 1 HVX client class */
+	HAP_POWER_STREAMING_2HVX_CLIENT_CLASS = 0x10,		/**< Camera streaming with 2 HVX client class */
+} HAP_power_app_type_payload;
+/**
+ * @} // HAP_power_enums
+ */
+
+/** Payload for HAP_power_set_linelock */
+typedef struct {
+	void* startAddress;					/**< Start address of the memory region to be locked. */
+	uint32 size;						/**< Size (bytes) of the memory region to be locked. Set size
+										*	to 0 to unlock memory. */
+	uint32 throttleBlockSize;			/**< Block size for throttling, in bytes;
+										* 0 for no throttling.  The region to be locked will be divided into
+										* blocks of this size for throttling purposes.
+										* Use for locking larger cache blocks.
+										* Applicable only when enabling line locking.Only ONE throttled linelock call is supported at this time.
+										* You can linelock additional regions (without throttling) using HAP_power_set_linelock_nothrottle*/
+	uint32 throttlePauseUs;				/**< Pause to be applied between locking each block, in microseconds. Applicable only when enabling line locking*/
+} HAP_power_linelock_payload;
+
+/** Payload for HAP_power_set_linelock_nothrottle */
+typedef struct {
+	void* startAddress;							/**< Start address of the memory region to be locked. */
+	uint32 size;								/**< Size (bytes) of the memory region to be locked. Set size to 0
+												* to unlock memory */
+} HAP_power_linelock_nothrottle_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+/** Option for dcvs payload */
+typedef enum {
+	HAP_DCVS_ADJUST_UP_DOWN =   0x1,		/**< increase and decrease core/bus clock speed. */
+	HAP_DCVS_ADJUST_ONLY_UP =   0x2,		/**< restricts DCVS from lowering the clock speed below the requested value . */
+} HAP_power_dcvs_payload_option;
+/**
+ * @} // HAP_power_enums
+ */
+
+/** Payload for HAP_power_set_DCVS */
+typedef struct {
+	boolean dcvs_enable;								/**< Set to TRUE to participate in DCVS, and FALSE otherwise. */
+	HAP_power_dcvs_payload_option dcvs_option;			/**< Set to one of
+														*		HAP_DCVS_ADJUST_UP_DOWN  - Allows for DCVS to adjust up and down.
+														*		HAP_DCVS_ADJUST_ONLY_UP  - Allows for DCVS to adjust up only. */
+} HAP_power_dcvs_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+/** Voltage corners for HAP DCVS V2 interface */
+typedef enum {
+	HAP_DCVS_VCORNER_DISABLE,
+	HAP_DCVS_VCORNER_SVS2,
+	HAP_DCVS_VCORNER_SVS,
+	HAP_DCVS_VCORNER_SVS_PLUS,
+	HAP_DCVS_VCORNER_NOM,
+	HAP_DCVS_VCORNER_NOM_PLUS,
+	HAP_DCVS_VCORNER_TURBO,
+	HAP_DCVS_VCORNER_TURBO_PLUS,
+	HAP_DCVS_VCORNER_MAX = 255,
+} HAP_dcvs_voltage_corner_t;
+/**
+ * @} // HAP_power_enums
+ */
+
+#define HAP_DCVS_VCORNER_SVSPLUS HAP_DCVS_VCORNER_SVS_PLUS
+#define HAP_DCVS_VCORNER_NOMPLUS HAP_DCVS_VCORNER_NOM_PLUS
+
+/** DCVS parameters for HAP_power_dcvs_v2_payload */
+typedef struct {
+	HAP_dcvs_voltage_corner_t target_corner;	/**< target voltage corner */
+	HAP_dcvs_voltage_corner_t min_corner;		/**< minimum voltage corner */
+	HAP_dcvs_voltage_corner_t max_corner;		/**< maximum voltage corner */
+	uint32 param1;								/**< reserved */
+	uint32 param2;								/**< reserved */
+	uint32 param3;								/**< reserved */
+} HAP_dcvs_params_t;
+
+/** Core clock parameters for HAP_power_dcvs_v3_payload */
+typedef struct {
+	HAP_dcvs_voltage_corner_t target_corner;	/**< target voltage corner */
+	HAP_dcvs_voltage_corner_t min_corner;		/**< minimum voltage corner */
+	HAP_dcvs_voltage_corner_t max_corner;		/**< maximum voltage corner */
+	uint32 param1;								/**< reserved */
+	uint32 param2;								/**< reserved */
+	uint32 param3;								/**< reserved */
+} HAP_core_params_t;
+
+/** Bus clock parameters for HAP_power_dcvs_v3_payload */
+typedef struct {
+	HAP_dcvs_voltage_corner_t target_corner;	/**< target voltage corner */
+	HAP_dcvs_voltage_corner_t min_corner;		/**< minimum voltage corner */
+	HAP_dcvs_voltage_corner_t max_corner;		/**< maximum voltage corner */
+	uint32 param1;								/**< reserved */
+	uint32 param2;								/**< reserved */
+	uint32 param3;								/**< reserved */
+} HAP_bus_params_t;
+
+/** DCVS v3 parameters for HAP_power_dcvs_v3_payload */
+typedef struct {
+	uint32 param1;					/**< reserved */
+	uint32 param2;					/**< reserved */
+	uint32 param3;					/**< reserved */
+	uint32 param4;					/**< reserved */
+	uint32 param5;					/**< reserved */
+	uint32 param6;					/**< reserved */
+} HAP_dcvs_v3_params_t;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+/** option for dcvs_v2 payload */
+typedef enum {
+	HAP_DCVS_V2_ADJUST_UP_DOWN =   0x1,					/**< Allows for DCVS to adjust up and down. */
+	HAP_DCVS_V2_ADJUST_ONLY_UP =   0x2,					/**< Allows for DCVS to adjust up only. */
+	HAP_DCVS_V2_POWER_SAVER_MODE = 0x4,					/**< HAP_DCVS_POWER_SAVER_MODE				-	Higher thresholds for power efficiency. */
+	HAP_DCVS_V2_POWER_SAVER_AGGRESSIVE_MODE = 0x8,		/**< HAP_DCVS_POWER_SAVER_AGGRESSIVE_MODE	-	Higher thresholds for power efficiency with faster ramp down. */
+	HAP_DCVS_V2_PERFORMANCE_MODE = 0x10,				/**< HAP_DCVS_PERFORMANCE_MODE				-	Lower thresholds for maximum performance */
+	HAP_DCVS_V2_DUTY_CYCLE_MODE = 0x20,					/**< HAP_DCVS_DUTY_CYCLE_MODE				-	only for HVX based clients.
+														*												For streaming class clients:
+														*													> detects periodicity based on HVX usage
+														*													> lowers clocks in the no HVX activity region of each period.
+														*												For compute class clients:
+														*													> Lowers clocks on no HVX activity detects and brings clocks up on detecting HVX activity again.
+														*													> Latency involved in bringing up the clock with be at max 1 to 2 ms. */
+
+
+
+} HAP_power_dcvs_v2_payload_option;
+/**
+ * @} // HAP_power_enums
+ */
+/** Payload for HAP_power_set_DCVS_v2 */
+typedef struct {
+	boolean dcvs_enable;								/**< Set to TRUE to participate in DCVS, and FALSE otherwise */
+	HAP_power_dcvs_v2_payload_option dcvs_option;		/**< Set to one of HAP_power_dcvs_v2_payload_option */
+	boolean set_latency;								/**< TRUE to set latency parameter, otherwise FALSE */
+	uint32 latency;										/**< sleep latency */
+	boolean set_dcvs_params;							/**< TRUE to set DCVS params, otherwise FALSE */
+	HAP_dcvs_params_t dcvs_params;						/**< DCVS parameters */
+} HAP_power_dcvs_v2_payload;
+
+/** Payload for HAP_power_set_DCVS_v3 */
+typedef struct {
+	boolean set_dcvs_enable;							/**< TRUE to consider DCVS enable/disable and option parameters, otherwise FALSE */
+	boolean dcvs_enable;								/**< Set to TRUE to participate in DCVS, and FALSE otherwise. */
+	HAP_power_dcvs_v2_payload_option dcvs_option;		/**< Set to one of HAP_power_dcvs_v2_payload_option */
+	boolean set_latency;								/**< TRUE to consider latency parameter, otherwise FALSE */
+	uint32 latency;										/**< sleep latency */
+	boolean set_core_params;							/**< TRUE to consider core clock params, otherwise FALSE */
+	HAP_core_params_t core_params;						/**< Core clock parameters */
+	boolean set_bus_params;								/**< TRUE to consider bus clock params, otherwise FALSE */
+	HAP_bus_params_t bus_params;						/**< Bus clock parameters */
+	boolean set_dcvs_v3_params;							/**< TRUE to consider DCVS v3 params, otherwise FALSE */
+	HAP_dcvs_v3_params_t dcvs_v3_params;				/**< DCVS v3 parameters */
+	boolean set_sleep_disable;							/**< TRUE to consider sleep disable/enable parameter, otherwise FALSE */
+	boolean sleep_disable;								/**< TRUE to disable sleep/LPM modes, FALSE to enable */
+} HAP_power_dcvs_v3_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+ /** Type for dcvs update request */
+typedef enum {
+	HAP_POWER_UPDATE_DCVS = 1,
+	HAP_POWER_UPDATE_SLEEP_LATENCY,
+	HAP_POWER_UPDATE_DCVS_PARAMS,
+} HAP_power_update_type_t;
+/**
+ * @} // HAP_power_enums
+ */
+/** Payload for DCVS update */
+typedef struct {
+	boolean dcvs_enable;							/**< TRUE for DCVS enable and FALSE for DCVS disable */
+	HAP_power_dcvs_v2_payload_option dcvs_option;	/**< Requested DCVS policy in case DCVS enable is TRUE */
+} HAP_power_update_dcvs_t;
+
+/** Payload for latency update */
+typedef struct {
+	boolean set_latency;							/**< TRUE if sleep latency request has to be considered */
+	unsigned int latency;							/**< Sleep latency request in micro seconds */
+} HAP_power_update_latency_t;
+
+/** Payload for DCVS params update */
+typedef struct {
+    boolean set_dcvs_params;						/**< Flag to mark DCVS params structure validity, TRUE for valid DCVS
+													*params request and FALSE otherwise */
+    HAP_dcvs_params_t dcvs_params;					/**< Intended DCVS params if set_dcvs_params is set to TRUE */
+} HAP_power_update_dcvs_params_t;
+
+/** Payload for HAP_power_set_DCVS_v2 */
+typedef struct {
+	HAP_power_update_type_t update_param;			/**< Type for which param to update */
+	union {
+		HAP_power_update_dcvs_t dcvs_payload;
+		HAP_power_update_latency_t latency_payload;
+		HAP_power_update_dcvs_params_t dcvs_params_payload;
+	};												/**< Update payload for DCVS, latency or DCVS params */
+} HAP_power_dcvs_v2_update_payload;
+
+/** Payload for HAP_power_set_streamer */
+typedef struct {
+	boolean set_streamer0_clk;				/**< Set streamer 0 clock */
+	boolean set_streamer1_clk;				/**< Set streamer 1 clock */
+	unsigned int streamer0_clkFreqHz;		/**< Streamer 0 clock frequency */
+	unsigned int streamer1_clkFreqHz;		/**< Streamer 1 clock frequency */
+	HAP_freq_match_type freqMatch;			/**< Clock frequency match */
+	uint32 param1;							/**< Reserved for future streamer parameters */
+	uint32 param2;							/**< Reserved for future streamer parameters */
+	uint32 param3;							/**< Reserved for future streamer parameters */
+} HAP_power_streamer_payload;
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+ /** Identifies the HAP power request type */
+typedef enum {
+	HAP_power_set_mips_bw = 1,				/**< Requests for MIPS. Provides
+											* fine-grained control to set MIPS values.
+											* Payload is set to HAP_power_payload */
+	HAP_power_set_HVX,						/**< Requests to enable / disable HVX
+											* Payload is set to HAP_power_hvx_payload */
+	HAP_power_set_apptype,					/**< Sets the app_type
+											* Payload is set to HAP_power_app_type_payload */
+	HAP_power_set_linelock,					/**< Sets the throttled L2 cache line locking parameters.
+											* Only one throttled call is supported at this time. Additional
+											* un-throttled line-locks can be performed using HAP_power_set_linelock_nothrottle
+											* Payload is set to HAP_power_linelock_payload */
+	HAP_power_set_DCVS,						/**< Requests to participate / stop participating in DCVS */
+	HAP_power_set_linelock_nothrottle,		/**< Sets the L2 cache line locking parameters (non-throttled).
+											* Payload is set to HAP_power_linelock_nothrottle_payload */
+	HAP_power_set_DCVS_v2,					/**< Requests to participate / stop participating in DCVS_v2 */
+	HAP_power_set_vapss,					/**< Sets the VAPSS core clock and DDR/IPNOC bandwidth
+											* Payload is set to HAP_power_vapss_payload */
+	HAP_power_set_vapss_v2,					/**< Sets the VAPSS core DMA/HCP clocks and DDR/IPNOC bandwidths
+											* Payload is set to HAP_power_vapss_payload_v2 */
+	HAP_power_set_dcvs_v2_update,			/**< Updates DCVS params
+											* Payload is set to HAP_power_dcvs_v2_update_payload */
+	HAP_power_set_streamer,					/**< Sets the streamer core clocks
+											* Payload is set to HAP_power_streamer_payload */
+	HAP_power_set_DCVS_v3,					/**< Updates DCVS params
+											* Payload is set to HAP_power_dcvs_v3_payload */
+	HAP_power_set_HMX,						/**< Requests to enable / disable HMX
+											* Payload is set to HAP_power_hmx_payload */
+} HAP_Power_request_type;
+/**
+ * @} // HAP_power_enums
+ */
+
+/** Data type to change power values on the DSP */
+typedef struct {
+	HAP_Power_request_type type;									/**< Identifies the request type */
+	union{
+		HAP_power_mips_bw_payload mips_bw;							/**< Requests for performance level */
+		HAP_power_vapss_payload vapss;								/**< Sets the VAPSS core clock and DDR/IPNOC bandwidth  */
+		HAP_power_vapss_payload_v2 vapss_v2;						/**< Sets the VAPSS core clock and DDR/IPNOC bandwidth  */
+		HAP_power_streamer_payload streamer;						/**< Sets the streamer core clocks */
+		HAP_power_hvx_payload hvx;									/**< Requests to enable / disable HVX */
+		HAP_power_app_type_payload apptype;							/**< Sets the app_type */
+		HAP_power_linelock_payload linelock;						/**< Sets the throttled L2 cache linelock parameters. Only one
+																	* throttled linelock is permitted at this time. Additional
+																	* un-throttled linelocks can be performed using linelock_nothrottle */
+		HAP_power_dcvs_payload dcvs;								/**< Updates DCVS params */
+		HAP_power_dcvs_v2_payload dcvs_v2;							/**< Updates DCVS_v2 params */
+		HAP_power_dcvs_v2_update_payload dcvs_v2_update;			/**< Updates DCVS_v2_update params */
+		HAP_power_linelock_nothrottle_payload linelock_nothrottle;	/**< Sets the un-throttled L2 cache linelock parameters */
+		HAP_power_dcvs_v3_payload dcvs_v3;							/**< Updates DCVS_v3 params */
+		HAP_power_hmx_payload hmx;									/**< Requests to turn on / off HMX */
+	};
+} HAP_power_request_t;
+
+/** @defgroup HAP_power_functions HAP POWER functions
+ *  @{
+ */
+/**
+* Method to set power values from the DSP
+* @param[in] context	-	To identify the power client
+* @param[in] request	-	Request params.
+* @retval 0 on success, AEE_EMMPMREGISTER on MMPM client register request failure, -1 on unknown error
+*/
+int HAP_power_set(void* context, HAP_power_request_t* request);
+/**
+ * @} // HAP_power_functions
+ */
+
+/** @defgroup HAP_power_enums HAP POWER enums
+ *  @{
+ */
+ /** Identifies the HAP power response type */
+typedef enum {
+	HAP_power_get_max_mips = 1,				/**< Returns the max mips supported (max_mips) */
+	HAP_power_get_max_bus_bw,				/**< Returns the max bus bandwidth supported (max_bus_bw) */
+	HAP_power_get_client_class,				/**< Returns the client class (client_class) */
+	HAP_power_get_clk_Freq,					/**< Returns the core clock frequency (clkFreqHz) */
+	HAP_power_get_aggregateAVSMpps,			/**< Returns the aggregate Mpps used by audio and voice (clkFreqHz) */
+	HAP_power_get_dcvsEnabled,				/**< Returns the dcvs status (enabled / disabled) */
+	HAP_power_get_vapss_core_clk_Freq,		/**< Returns the VAPSS core clock frequency (clkFreqHz) */
+	HAP_power_get_dma_core_clk_Freq,		/**< Returns the DMA core clock frequency (clkFreqHz) */
+	HAP_power_get_hcp_core_clk_Freq,		/**< Returns the HCP core clock frequency (clkFreqHz) */
+	HAP_power_get_streamer0_core_clk_Freq,	/**< Returns the streamer 0 core clock frequency (clkFreqHz) */
+	HAP_power_get_streamer1_core_clk_Freq,	/**< Returns the streamer 1 core clock frequency (clkFreqHz) */
+} HAP_Power_response_type;
+/**
+ * @} // HAP_power_enums
+ */
+
+/** Data type to retrieve power values from the DSP */
+typedef struct {
+	HAP_Power_response_type type;			/**< Identifies the type to retrieve. */
+	union{
+		unsigned int max_mips;				/**< Max mips supported */
+		uint64 max_bus_bw;					/**< Max bus bw supported */
+		unsigned int client_class;			/**< Current client class */
+		unsigned int clkFreqHz;				/**< Current core CPU frequency */
+		unsigned int aggregateAVSMpps;		/**< Aggregate AVS Mpps used by audio and voice */
+		boolean dcvsEnabled;				/**< Indicates if dcvs is enabled / disabled. */
+	};
+} HAP_power_response_t;
+
+/** @defgroup HAP_power_functions HAP POWER functions
+ *  @{
+ */
+
+/**
+* Method to retrieve power values from the DSP
+* @param[in] context	-	Ignored
+* @param[out] response	-	Response.
+*/
+int HAP_power_get(void* context, HAP_power_response_t* response);
+
+/**
+* Method to initialize dcvs v3 structure in request param. It enables
+*		flags and resets params for all fields in dcvs v3. So, this
+*		can also be used to remove applied dcvs v3 params and restore
+*		defaults.
+* @param[in] request	-	Pointer to request params.
+*/
+/*static inline void HAP_power_set_dcvs_v3_init(HAP_power_request_t* request) {
+	memset(request, 0, sizeof(HAP_power_request_t) );
+	request->type = HAP_power_set_DCVS_v3;
+	request->dcvs_v3.set_dcvs_enable = TRUE;
+	request->dcvs_v3.dcvs_enable = TRUE;
+	request->dcvs_v3.dcvs_option = HAP_DCVS_V2_POWER_SAVER_MODE;
+	request->dcvs_v3.set_latency = TRUE;
+	request->dcvs_v3.latency = 65535;
+	request->dcvs_v3.set_core_params = TRUE;
+	request->dcvs_v3.set_bus_params = TRUE;
+	request->dcvs_v3.set_dcvs_v3_params = TRUE;
+	request->dcvs_v3.set_sleep_disable = TRUE;
+	return;
+}*/
+
+/**
+* Method to enable/disable dcvs and set particular dcvs policy.
+* @param[in] context		-	User context.
+* @param[in] dcvs_enable	-	TRUE to enable dcvs, FALSE to disable dcvs.
+* @param[in] dcvs_option	-	To set particular dcvs policy. In case of dcvs disable
+*                           request, this param will be ignored.
+* @returns	-	0 on success
+*/
+/*static inline int HAP_power_set_dcvs_option(void* context, boolean dcvs_enable,
+		HAP_power_dcvs_v2_payload_option dcvs_option) {
+	HAP_power_request_t request;
+	memset(&request, 0, sizeof(HAP_power_request_t) );
+	request.type = HAP_power_set_DCVS_v3;
+	request.dcvs_v3.set_dcvs_enable = TRUE;
+	request.dcvs_v3.dcvs_enable = dcvs_enable;
+	if(dcvs_enable)
+		request.dcvs_v3.dcvs_option = dcvs_option;
+	return HAP_power_set(context, &request);
+}*/
+
+/**
+* Method to set/reset sleep latency.
+* @param[in] context	-	User context.
+* @param[in] latency	-	Sleep latency value in microseconds, should be > 1.
+*						Use 65535 max value to reset it to default.
+* @returns	-	0 on success
+*/
+/*static inline int HAP_power_set_sleep_latency(void* context, uint32 latency) {
+	HAP_power_request_t request;
+	memset(&request, 0, sizeof(HAP_power_request_t) );
+	request.type = HAP_power_set_DCVS_v3;
+	request.dcvs_v3.set_latency = TRUE;
+	request.dcvs_v3.latency = latency;
+	return HAP_power_set(context, &request);
+}*/
+
+/**
+* Method to set/reset DSP core clock voltage corners.
+* @param[in] context		-	User context.
+* @param[in] target_corner	-	Target voltage corner.
+* @param[in] min_corner		-	Minimum voltage corner.
+* @param[in] max_corner		-	Maximum voltage corner.
+* @returns	-	0 on success
+*/
+/*static inline int HAP_power_set_core_corner(void* context, uint32 target_corner,
+		uint32 min_corner, uint32 max_corner) {
+	HAP_power_request_t request;
+	memset(&request, 0, sizeof(HAP_power_request_t) );
+	request.type = HAP_power_set_DCVS_v3;
+	request.dcvs_v3.set_core_params = TRUE;
+	request.dcvs_v3.core_params.min_corner = (HAP_dcvs_voltage_corner_t) (min_corner);
+	request.dcvs_v3.core_params.max_corner = (HAP_dcvs_voltage_corner_t) (max_corner);
+	request.dcvs_v3.core_params.target_corner = (HAP_dcvs_voltage_corner_t) (target_corner);
+	return HAP_power_set(context, &request);
+}*/
+
+/**
+* Method to set/reset bus clock voltage corners.
+* @param[in] context		-	User context.
+* @param[in] target_corner	-	Target voltage corner.
+* @param[in] min_corner		-	Minimum voltage corner.
+* @param[in] max_corner		-	Maximum voltage corner.
+* @returns	-	0 on success
+*/
+/*static inline int HAP_power_set_bus_corner(void* context, uint32 target_corner,
+		uint32 min_corner, uint32 max_corner) {
+    HAP_power_request_t request;
+	memset(&request, 0, sizeof(HAP_power_request_t) );
+	request.type = HAP_power_set_DCVS_v3;
+	request.dcvs_v3.set_bus_params = TRUE;
+	request.dcvs_v3.bus_params.min_corner = (HAP_dcvs_voltage_corner_t) (min_corner);
+	request.dcvs_v3.bus_params.max_corner = (HAP_dcvs_voltage_corner_t) (max_corner);
+	request.dcvs_v3.bus_params.target_corner = (HAP_dcvs_voltage_corner_t) (target_corner);
+	return HAP_power_set(context, &request);
+}*/
+
+/**
+* Method to disable/enable all low power modes.
+* @param[in] context		-	User context.
+* @param[in] sleep_disable	-	TRUE to disable all low power modes.
+*							FALSE to re-enable all low power modes.
+* @returns	-	0 on success
+*/
+/*static inline int HAP_power_set_sleep_mode(void* context, boolean sleep_disable) {
+	HAP_power_request_t request;
+	memset(&request, 0, sizeof(HAP_power_request_t) );
+	request.type = HAP_power_set_DCVS_v3;
+	request.dcvs_v3.set_sleep_disable = TRUE;
+	request.dcvs_v3.sleep_disable = sleep_disable;
+	return HAP_power_set(context, &request);
+}*/
+
+
+/**
+* This API is deprecated and might generate undesired results.
+* Please use the HAP_power_get() and HAP_power_set() APIs instead.
+* Requests a performance level by percentage for clock speed
+* and bus speed.  Passing 0 for any parameter results in no
+* request being issued for that particular attribute.
+* @param[in] clock		-	percentage of target's maximum clock speed
+* @param[in] bus		-	percentage of target's maximum bus speed
+* @param[in] latency	-	maximum hardware wake up latency in microseconds.  The
+*						higher the value the deeper state of sleep
+*						that can be entered but the longer it may
+*						take to awaken.
+* @retval 0 on success
+* @par Comments	:	Performance metrics vary from target to target so the
+*					intent of this API is to allow callers to set a relative
+*					performance level to achieve the desired balance between
+*					performance and power saving.
+*/
+int HAP_power_request(int clock, int bus, int latency);
+
+/**
+* This API is deprecated and might generate undesired results.
+* Please use the HAP_power_get() and HAP_power_set() APIs instead.
+* Requests a performance level by absolute values.  Passing 0
+* for any parameter results in no request being issued for that
+* particular attribute.
+* @param[in] clock		-	speed in MHz
+* @param[in] bus		-	bus speed in MHz
+* @param[in] latency	-	maximum hardware wakeup latency in microseconds.  The
+*						higher the value the deeper state of
+*						sleep that can be entered but the
+*						longer it may take to awaken.
+* @retval 0 on success
+* @par Comments	:	This API allows callers who are aware of their target
+*					specific capabilities to set them explicitly.
+*/
+int HAP_power_request_abs(int clock, int bus, int latency);
+
+/**
+* This API is deprecated and might generate undesired results.
+* Please use the HAP_power_get() and HAP_power_set() APIs instead.
+* queries the target for its clock and bus speed capabilities
+* @param[out] clock_max	-	maximum clock speed supported in MHz
+* @param[out] bus_max	-	maximum bus speed supported in MHz
+* @retval 0 on success
+*/
+int HAP_power_get_max_speed(int* clock_max, int* bus_max);
+
+/**
+* This API is deprecated and might generate undesired results.
+* Please use the HAP_power_get() and HAP_power_set() APIs instead.
+* Upvote for HVX power
+* @retval 0 on success
+*/
+int HVX_power_request(void);
+
+/**
+* This API is deprecated and might generate undesired results.
+* Please use the HAP_power_get() and HAP_power_set() APIs instead.
+* Downvote for HVX power
+* @retval 0 on success
+*/
+int HVX_power_release(void);
+
+/**
+* Method to destroy clients created through HAP_power_set
+* @param[in] context	-	To uniquely identify the client
+* @retval 0 on success, AEE_ENOSUCHCLIENT on Invalid context, -1 on unknown error
+* @brief DO NOT call this API directly, use HAP_power_destroy instead.
+*/
+int HAP_power_destroy_client(void *context);
+
+/**
+* @param[in] client	-	To uniquely identify the client context.
+* @retval 0 on success, AEE_EUNSUPPORTEDAPI if the API is not supported on the DSP image, AEE_ENOSUCHCLIENT on Invalid context, -1 on unknown error
+* @brief Method to destroy clients created through HAP_power_set, wrapper to HAP_power_destroy_client API
+*/
+static inline int HAP_power_destroy(void *client){
+	if(0 != HAP_power_destroy_client)
+		return HAP_power_destroy_client(client);
+	return AEE_EUNSUPPORTEDAPI;
+}
+
+/**
+* Method to create user client context
+* @retval context for client
+*/
+//static inline void* HAP_utils_create_context(void) {
+	/*
+	 * Allocate 1 byte of memory for a unique context identifier
+	 * Clients can also allocate memory and use it as unique context identifier
+	 */
+//	return malloc(1);
+//}
+
+/**
+* Method to destroy user client context
+* @param context of client
+*/
+/*static inline void HAP_utils_destroy_context(void* context) {
+	free(context);
+}*/
+
+/**
+ * @} // HAP_power_functions
+ */
+#ifdef __cplusplus
+}
+#endif
+#endif //_HAP_POWER_H
+
@@ -0,0 +1,319 @@
+/*
+ * Copyright (c) 2012-2018, The Linux Foundation. All rights reserved.
+ *
+ * This program is free software; you can redistribute it and/or modify
+ * it under the terms of the GNU General Public License version 2 and
+ * only version 2 as published by the Free Software Foundation.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ */
+#ifndef ADSPRPC_SHARED_H
+#define ADSPRPC_SHARED_H
+
+#include <stdint.h>
+#include <stddef.h>
+#include <sys/types.h>
+#include <linux/types.h>
+
+#define FASTRPC_IOCTL_INVOKE	_IOWR('R', 1, struct fastrpc_ioctl_invoke)
+#define FASTRPC_IOCTL_MMAP	_IOWR('R', 2, struct fastrpc_ioctl_mmap)
+#define FASTRPC_IOCTL_MUNMAP	_IOWR('R', 3, struct fastrpc_ioctl_munmap)
+#define FASTRPC_IOCTL_MMAP_64	_IOWR('R', 14, struct fastrpc_ioctl_mmap_64)
+#define FASTRPC_IOCTL_MUNMAP_64	_IOWR('R', 15, struct fastrpc_ioctl_munmap_64)
+#define FASTRPC_IOCTL_INVOKE_FD	_IOWR('R', 4, struct fastrpc_ioctl_invoke_fd)
+#define FASTRPC_IOCTL_SETMODE	_IOWR('R', 5, uint32_t)
+#define FASTRPC_IOCTL_INIT	_IOWR('R', 6, struct fastrpc_ioctl_init)
+#define FASTRPC_IOCTL_INVOKE_ATTRS \
+				_IOWR('R', 7, struct fastrpc_ioctl_invoke_attrs)
+#define FASTRPC_IOCTL_GETINFO	_IOWR('R', 8, uint32_t)
+#define FASTRPC_IOCTL_GETPERF	_IOWR('R', 9, struct fastrpc_ioctl_perf)
+#define FASTRPC_IOCTL_INIT_ATTRS _IOWR('R', 10, struct fastrpc_ioctl_init_attrs)
+#define FASTRPC_IOCTL_INVOKE_CRC _IOWR('R', 11, struct fastrpc_ioctl_invoke_crc)
+#define FASTRPC_IOCTL_CONTROL   _IOWR('R', 12, struct fastrpc_ioctl_control)
+#define FASTRPC_IOCTL_MUNMAP_FD _IOWR('R', 13, struct fastrpc_ioctl_munmap_fd)
+
+#define FASTRPC_GLINK_GUID "fastrpcglink-apps-dsp"
+#define FASTRPC_SMD_GUID "fastrpcsmd-apps-dsp"
+#define DEVICE_NAME      "adsprpc-smd"
+
+/* Set for buffers that have no virtual mapping in userspace */
+#define FASTRPC_ATTR_NOVA 0x1
+
+/* Set for buffers that are NOT dma coherent */
+#define FASTRPC_ATTR_NON_COHERENT 0x2
+
+/* Set for buffers that are dma coherent */
+#define FASTRPC_ATTR_COHERENT 0x4
+
+/* Fastrpc attribute for keeping the map persistent */
+#define FASTRPC_ATTR_KEEP_MAP	0x8
+
+/* Fastrpc attribute for no map */
+#define FASTRPC_ATTR_NOMAP   (16)
+
+/* Driver should operate in parallel with the co-processor */
+#define FASTRPC_MODE_PARALLEL    0
+
+/* Driver should operate in serial mode with the co-processor */
+#define FASTRPC_MODE_SERIAL      1
+
+/* Driver should operate in profile mode with the co-processor */
+#define FASTRPC_MODE_PROFILE     2
+
+/* Set FastRPC session ID to 1 */
+#define FASTRPC_MODE_SESSION     4
+
+/* INIT a new process or attach to guestos */
+#define FASTRPC_INIT_ATTACH      0
+#define FASTRPC_INIT_CREATE      1
+#define FASTRPC_INIT_CREATE_STATIC  2
+#define FASTRPC_INIT_ATTACH_SENSORS 3
+
+/* Retrives number of input buffers from the scalars parameter */
+#define REMOTE_SCALARS_INBUFS(sc)        (((sc) >> 16) & 0x0ff)
+
+/* Retrives number of output buffers from the scalars parameter */
+#define REMOTE_SCALARS_OUTBUFS(sc)       (((sc) >> 8) & 0x0ff)
+
+/* Retrives number of input handles from the scalars parameter */
+#define REMOTE_SCALARS_INHANDLES(sc)     (((sc) >> 4) & 0x0f)
+
+/* Retrives number of output handles from the scalars parameter */
+#define REMOTE_SCALARS_OUTHANDLES(sc)    ((sc) & 0x0f)
+
+#define REMOTE_SCALARS_LENGTH(sc)	(REMOTE_SCALARS_INBUFS(sc) +\
+					REMOTE_SCALARS_OUTBUFS(sc) +\
+					REMOTE_SCALARS_INHANDLES(sc) +\
+					REMOTE_SCALARS_OUTHANDLES(sc))
+
+#define REMOTE_SCALARS_MAKEX(attr, method, in, out, oin, oout) \
+		((((uint32_t)   (attr) & 0x7) << 29) | \
+		(((uint32_t) (method) & 0x1f) << 24) | \
+		(((uint32_t)     (in) & 0xff) << 16) | \
+		(((uint32_t)    (out) & 0xff) <<  8) | \
+		(((uint32_t)    (oin) & 0x0f) <<  4) | \
+		((uint32_t)   (oout) & 0x0f))
+
+#define REMOTE_SCALARS_MAKE(method, in, out) \
+		REMOTE_SCALARS_MAKEX(0, method, in, out, 0, 0)
+
+
+#ifndef VERIFY_PRINT_ERROR
+#define VERIFY_EPRINTF(format, args) (void)0
+#endif
+
+#ifndef VERIFY_PRINT_INFO
+#define VERIFY_IPRINTF(args) (void)0
+#endif
+
+#ifndef VERIFY
+#define __STR__(x) #x ":"
+#define __TOSTR__(x) __STR__(x)
+#define __FILE_LINE__ __FILE__ ":" __TOSTR__(__LINE__)
+
+#define VERIFY(err, val) \
+do {\
+	VERIFY_IPRINTF(__FILE_LINE__"info: calling: " #val "\n");\
+	if ((val) == 0) {\
+		(err) = (err) == 0 ? -1 : (err);\
+		VERIFY_EPRINTF(__FILE_LINE__"error: %d: " #val "\n", (err));\
+	} else {\
+		VERIFY_IPRINTF(__FILE_LINE__"info: passed: " #val "\n");\
+	} \
+} while (0)
+#endif
+
+#define remote_arg64_t    union remote_arg64
+
+struct remote_buf64 {
+	uint64_t pv;
+	uint64_t len;
+};
+
+struct remote_dma_handle64 {
+	int fd;
+	uint32_t offset;
+	uint32_t len;
+};
+
+union remote_arg64 {
+	struct remote_buf64	buf;
+	struct remote_dma_handle64 dma;
+	uint32_t h;
+};
+
+#define remote_arg_t    union remote_arg
+
+struct remote_buf {
+	void *pv;		/* buffer pointer */
+	size_t len;		/* length of buffer */
+};
+
+struct remote_dma_handle {
+	int fd;
+	uint32_t offset;
+};
+
+union remote_arg {
+	struct remote_buf buf;	/* buffer info */
+	struct remote_dma_handle dma;
+	uint32_t h;		/* remote handle */
+};
+
+struct fastrpc_ioctl_invoke {
+	uint32_t handle;	/* remote handle */
+	uint32_t sc;		/* scalars describing the data */
+	remote_arg_t *pra;	/* remote arguments list */
+};
+
+struct fastrpc_ioctl_invoke_fd {
+	struct fastrpc_ioctl_invoke inv;
+	int *fds;		/* fd list */
+};
+
+struct fastrpc_ioctl_invoke_attrs {
+	struct fastrpc_ioctl_invoke inv;
+	int *fds;		/* fd list */
+	unsigned int *attrs;	/* attribute list */
+};
+
+struct fastrpc_ioctl_invoke_crc {
+	struct fastrpc_ioctl_invoke inv;
+	int *fds;		/* fd list */
+	unsigned int *attrs;	/* attribute list */
+	unsigned int *crc;
+};
+
+struct fastrpc_ioctl_init {
+	uint32_t flags;		/* one of FASTRPC_INIT_* macros */
+	uintptr_t file;		/* pointer to elf file */
+	uint32_t filelen;	/* elf file length */
+	int32_t filefd;		/* ION fd for the file */
+	uintptr_t mem;		/* mem for the PD */
+	uint32_t memlen;	/* mem length */
+	int32_t memfd;		/* ION fd for the mem */
+};
+
+struct fastrpc_ioctl_init_attrs {
+		struct fastrpc_ioctl_init init;
+		int attrs;
+		unsigned int siglen;
+};
+
+struct fastrpc_ioctl_munmap {
+	uintptr_t vaddrout;	/* address to unmap */
+	size_t size;		/* size */
+};
+
+struct fastrpc_ioctl_munmap_64 {
+	uint64_t vaddrout;	/* address to unmap */
+	size_t size;		/* size */
+};
+
+struct fastrpc_ioctl_mmap {
+	int fd;				/* ion fd */
+	uint32_t flags;			/* flags for dsp to map with */
+	uintptr_t vaddrin;		/* optional virtual address */
+	size_t size;			/* size */
+	uintptr_t vaddrout;		/* dsps virtual address */
+};
+
+struct fastrpc_ioctl_mmap_64 {
+	int fd;					/* ion fd */
+	uint32_t flags;			/* flags for dsp to map with */
+	uint64_t vaddrin;		/* optional virtual address */
+	size_t size;			/* size */
+	uint64_t vaddrout;		/* dsps virtual address */
+};
+
+struct fastrpc_ioctl_munmap_fd {
+	int     fd;				/* fd */
+	uint32_t  flags;		/* control flags */
+	uintptr_t va;			/* va */
+	ssize_t  len;			/* length */
+};
+
+struct fastrpc_ioctl_perf {			/* kernel performance data */
+	uintptr_t data;
+	uint32_t numkeys;
+	uintptr_t keys;
+};
+
+#define FASTRPC_CONTROL_LATENCY (1)
+struct fastrpc_ctrl_latency {
+	uint32_t enable;		/* latency control enable */
+	uint32_t level;			/* level of control */
+};
+
+#define FASTRPC_CONTROL_SMMU (2)
+struct fastrpc_ctrl_smmu {
+	uint32_t sharedcb;
+};
+
+#define FASTRPC_CONTROL_KALLOC (3)
+struct fastrpc_ctrl_kalloc {
+	uint32_t kalloc_support; /* Remote memory allocation from kernel */
+};
+
+struct fastrpc_ioctl_control {
+	uint32_t req;
+	union {
+		struct fastrpc_ctrl_latency lp;
+		struct fastrpc_ctrl_smmu smmu;
+		struct fastrpc_ctrl_kalloc kalloc;
+	};
+};
+
+struct smq_null_invoke {
+	uint64_t ctx;			/* invoke caller context */
+	uint32_t handle;	    /* handle to invoke */
+	uint32_t sc;		    /* scalars structure describing the data */
+};
+
+struct smq_phy_page {
+	uint64_t addr;		/* physical address */
+	uint64_t size;		/* size of contiguous region */
+};
+
+struct smq_invoke_buf {
+	int num;		/* number of contiguous regions */
+	int pgidx;		/* index to start of contiguous region */
+};
+
+struct smq_invoke {
+	struct smq_null_invoke header;
+	struct smq_phy_page page;   /* remote arg and list of pages address */
+};
+
+struct smq_msg {
+	uint32_t pid;           /* process group id */
+	uint32_t tid;           /* thread id */
+	struct smq_invoke invoke;
+};
+
+struct smq_invoke_rsp {
+	uint64_t ctx;			/* invoke caller context */
+	int retval;	             /* invoke return value */
+};
+
+static inline struct smq_invoke_buf *smq_invoke_buf_start(remote_arg64_t *pra,
+							uint32_t sc)
+{
+	unsigned int len = REMOTE_SCALARS_LENGTH(sc);
+
+	return (struct smq_invoke_buf *)(&pra[len]);
+}
+
+static inline struct smq_phy_page *smq_phy_page_start(uint32_t sc,
+						struct smq_invoke_buf *buf)
+{
+	unsigned int nTotal = REMOTE_SCALARS_LENGTH(sc);
+
+	return (struct smq_phy_page *)(&buf[nTotal]);
+}
+
+#endif
@@ -0,0 +1,208 @@
+/**
+ * Copyright (c) 2019, The Linux Foundation. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are
+ * met:
+ *    * Redistributions of source code must retain the above copyright
+ *      notice, this list of conditions and the following disclaimer.
+ *    * Redistributions in binary form must reproduce the above
+ *      copyright notice, this list of conditions and the following
+ *      disclaimer in the documentation and/or other materials provided
+ *      with the distribution.
+ *    * Neither the name of The Linux Foundation nor the names of its
+ *      contributors may be used to endorse or promote products derived
+ *      from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED "AS IS" AND ANY EXPRESS OR IMPLIED
+ * WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NON-INFRINGEMENT
+ * ARE DISCLAIMED.  IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS
+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR
+ * BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY,
+ * WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
+ * OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN
+ * IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#ifndef _APPS_STD_H
+#define _APPS_STD_H
+#include "AEEStdDef.h"
+#ifndef __QAIC_HEADER
+#define __QAIC_HEADER(ff) ff
+#endif //__QAIC_HEADER
+
+#ifndef __QAIC_HEADER_EXPORT
+#define __QAIC_HEADER_EXPORT
+#endif // __QAIC_HEADER_EXPORT
+
+#ifndef __QAIC_HEADER_ATTRIBUTE
+#define __QAIC_HEADER_ATTRIBUTE
+#endif // __QAIC_HEADER_ATTRIBUTE
+
+#ifndef __QAIC_IMPL
+#define __QAIC_IMPL(ff) ff
+#endif //__QAIC_IMPL
+
+#ifndef __QAIC_IMPL_EXPORT
+#define __QAIC_IMPL_EXPORT
+#endif // __QAIC_IMPL_EXPORT
+
+#ifndef __QAIC_IMPL_ATTRIBUTE
+#define __QAIC_IMPL_ATTRIBUTE
+#endif // __QAIC_IMPL_ATTRIBUTE
+#ifdef __cplusplus
+extern "C" {
+#endif
+#if !defined(__QAIC_STRING1_OBJECT_DEFINED__) && !defined(__STRING1_OBJECT__)
+#define __QAIC_STRING1_OBJECT_DEFINED__
+#define __STRING1_OBJECT__
+typedef struct _cstring1_s {
+   char* data;
+   int dataLen;
+} _cstring1_t;
+
+#endif /* __QAIC_STRING1_OBJECT_DEFINED__ */
+/**
+    * standard library functions remoted from the apps to the dsp
+    */
+typedef int apps_std_FILE;
+enum apps_std_SEEK {
+   APPS_STD_SEEK_SET,
+   APPS_STD_SEEK_CUR,
+   APPS_STD_SEEK_END,
+   _32BIT_PLACEHOLDER_apps_std_SEEK = 0x7fffffff
+};
+typedef enum apps_std_SEEK apps_std_SEEK;
+typedef struct apps_std_DIR apps_std_DIR;
+struct apps_std_DIR {
+   uint64 handle;
+};
+typedef struct apps_std_DIRENT apps_std_DIRENT;
+struct apps_std_DIRENT {
+   int ino;
+   char name[255];
+};
+typedef struct apps_std_STAT apps_std_STAT;
+struct apps_std_STAT {
+   uint64 tsz;
+   uint64 dev;
+   uint64 ino;
+   uint32 mode;
+   uint32 nlink;
+   uint64 rdev;
+   uint64 size;
+   int64 atime;
+   int64 atimensec;
+   int64 mtime;
+   int64 mtimensec;
+   int64 ctime;
+   int64 ctimensec;
+};
+/**
+       * @retval, if operation fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fopen)(const char* name, const char* mode, apps_std_FILE* psout) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_freopen)(apps_std_FILE sin, const char* name, const char* mode, apps_std_FILE* psout) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fflush)(apps_std_FILE sin) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fclose)(apps_std_FILE sin) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @param,  bEOF, if read or write bytes <= bufLen bytes then feof() is called
+       *          and the result is returned in bEOF, otherwise bEOF is set to 0.
+       * @retval, if read or write return 0 for non zero length buffers, ferror is checked
+       *          and a non zero value is returned in case of error with no rout parameters
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fread)(apps_std_FILE sin, byte* buf, int bufLen, int* bytesRead, int* bEOF) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fwrite)(apps_std_FILE sin, const byte* buf, int bufLen, int* bytesWritten, int* bEOF) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @param, pos, this buffer is filled up to MIN(posLen, sizeof(fpos_t))
+       * @param, posLenReq, returns sizeof(fpos_t)
+       * @retval, if operation fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fgetpos)(apps_std_FILE sin, byte* pos, int posLen, int* posLenReq) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @param, if size of pos doesn't match the system size an error is returned.
+       *         fgetpos can be used to query the size of fpos_t
+       * @retval, if operation fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fsetpos)(apps_std_FILE sin, const byte* pos, int posLen) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @retval, if operation fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_ftell)(apps_std_FILE sin, int* pos) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fseek)(apps_std_FILE sin, int offset, apps_std_SEEK whence) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_flen)(apps_std_FILE sin, uint64* len) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @retval, only fails if transport fails
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_rewind)(apps_std_FILE sin) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_feof)(apps_std_FILE sin, int* bEOF) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_ferror)(apps_std_FILE sin, int* err) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_clearerr)(apps_std_FILE sin) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_print_string)(const char* str) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @param val, must contain space for NULL
+       * @param valLenReq, length required with NULL
+       * @retval, if fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_getenv)(const char* name, char* val, int valLen, int* valLenReq) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * @retval, if fails errno is returned
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_setenv)(const char* name, const char* val, int override) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_unsetenv)(const char* name) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * This function will try to open a file given directories in envvarname separated by
+       * delim.
+       * so given environment variable FOO_PATH=/foo;/bar
+       *    fopen_wth_env("FOO_PATH", ";", "path/to/file", "rw", &out);
+       * will try to open /foo/path/to/file, /bar/path/to/file
+       * if the variable is unset, it will open the file directly
+       *
+       * @param envvarname, name of the environment variable containing the path
+       * @param delim, delimiator string, such as ";"
+       * @param name, name of the file
+       * @param mode, mode
+       * @param psout, output handle
+       * @retval, 0 on success errno or -1 on failure
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fopen_with_env)(const char* envvarname, const char* delim, const char* name, const char* mode, apps_std_FILE* psout) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fgets)(apps_std_FILE sin, byte* buf, int bufLen, int* bEOF) __QAIC_HEADER_ATTRIBUTE;
+/**
+       * This method will return the paths that are searched when looking for a file.
+       *  The paths are defined by the environment variable (separated by delimiters)
+       *  that is passed to the method.
+       *
+       * @param envvarname, name of the environment variable containing the path
+       * @param delim, delimiator string, such as ";"
+       * @param name, name of the file
+       * @param paths, Search paths
+       * @param numPaths, Actual number of paths found
+       * @param maxPathLen, The max path length
+       * @retval, 0 on success errno or -1 on failure
+       *
+       */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_get_search_paths_with_env)(const char* envvarname, const char* delim, _cstring1_t* paths, int pathsLen, uint32* numPaths, uint16* maxPathLen) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fileExists)(const char* path, boolean* exists) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fsync)(apps_std_FILE sin) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fremove)(const char* name) __QAIC_HEADER_ATTRIBUTE;
+/**
+        * This function decrypts the file using the provided open file descriptor, closes the
+        * original descriptor and return a new file descriptor.
+        * @retval, if operation fails errno is returned
+        */
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_fdopen_decrypt)(apps_std_FILE sin, apps_std_FILE* psout) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_opendir)(const char* name, apps_std_DIR* dir) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_closedir)(const apps_std_DIR* dir) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_readdir)(const apps_std_DIR* dir, apps_std_DIRENT* dirent, int* bEOF) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_mkdir)(const char* name, int mode) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_rmdir)(const char* name) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_stat)(const char* name, apps_std_STAT* stat) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_ftrunc)(apps_std_FILE sin, int offset) __QAIC_HEADER_ATTRIBUTE;
+__QAIC_HEADER_EXPORT int __QAIC_HEADER(apps_std_frename)(const char* oldname, const char* newname) __QAIC_HEADER_ATTRIBUTE;
+#ifdef __cplusplus
+}
+#endif
+#endif //_APPS_STD_H
@@ -0,0 +1,204 @@
+/*
+ * drivers/staging/android/uapi/ion.h
+ *
+ * Copyright (C) 2011 Google, Inc.
+ *
+ * This software is licensed under the terms of the GNU General Public
+ * License version 2, as published by the Free Software Foundation, and
+ * may be copied, distributed, and modified under those terms.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ */
+
+#ifndef _UAPI_LINUX_ION_H
+#define _UAPI_LINUX_ION_H
+
+#include <stddef.h>
+#include <linux/ioctl.h>
+#include <linux/types.h>
+
+typedef int ion_user_handle_t;
+
+/**
+ * enum ion_heap_types - list of all possible types of heaps
+ * @ION_HEAP_TYPE_SYSTEM:	 memory allocated via vmalloc
+ * @ION_HEAP_TYPE_SYSTEM_CONTIG: memory allocated via kmalloc
+ * @ION_HEAP_TYPE_CARVEOUT:	 memory allocated from a prereserved
+ *				 carveout heap, allocations are physically
+ *				 contiguous
+ * @ION_HEAP_TYPE_DMA:		 memory allocated via DMA API
+ * @ION_NUM_HEAPS:		 helper for iterating over heaps, a bit mask
+ *				 is used to identify the heaps, so only 32
+ *				 total heap types are supported
+ */
+enum ion_heap_type {
+	ION_HEAP_TYPE_SYSTEM,
+	ION_HEAP_TYPE_SYSTEM_CONTIG,
+	ION_HEAP_TYPE_CARVEOUT,
+	ION_HEAP_TYPE_CHUNK,
+	ION_HEAP_TYPE_DMA,
+	ION_HEAP_TYPE_CUSTOM, /*
+			       * must be last so device specific heaps always
+			       * are at the end of this enum
+			       */
+	ION_NUM_HEAPS = 16,
+};
+
+#define ION_HEAP_SYSTEM_MASK		((1 << ION_HEAP_TYPE_SYSTEM))
+#define ION_HEAP_SYSTEM_CONTIG_MASK	((1 << ION_HEAP_TYPE_SYSTEM_CONTIG))
+#define ION_HEAP_CARVEOUT_MASK		((1 << ION_HEAP_TYPE_CARVEOUT))
+#define ION_HEAP_TYPE_DMA_MASK		((1 << ION_HEAP_TYPE_DMA))
+
+#define ION_NUM_HEAP_IDS		(sizeof(unsigned int) * 8)
+
+/**
+ * allocation flags - the lower 16 bits are used by core ion, the upper 16
+ * bits are reserved for use by the heaps themselves.
+ */
+#define ION_FLAG_CACHED 1		/*
+					 * mappings of this buffer should be
+					 * cached, ion will do cache
+					 * maintenance when the buffer is
+					 * mapped for dma
+					 */
+#define ION_FLAG_CACHED_NEEDS_SYNC 2	/*
+					 * mappings of this buffer will created
+					 * at mmap time, if this is set
+					 * caches must be managed
+					 * manually
+					 */
+
+/**
+ * DOC: Ion Userspace API
+ *
+ * create a client by opening /dev/ion
+ * most operations handled via following ioctls
+ *
+ */
+
+/**
+ * struct ion_allocation_data - metadata passed from userspace for allocations
+ * @len:		size of the allocation
+ * @align:		required alignment of the allocation
+ * @heap_id_mask:	mask of heap ids to allocate from
+ * @flags:		flags passed to heap
+ * @handle:		pointer that will be populated with a cookie to use to
+ *			refer to this allocation
+ *
+ * Provided by userspace as an argument to the ioctl
+ */
+struct ion_allocation_data {
+	size_t len;
+	size_t align;
+	unsigned int heap_id_mask;
+	unsigned int flags;
+	ion_user_handle_t handle;
+};
+
+/**
+ * struct ion_fd_data - metadata passed to/from userspace for a handle/fd pair
+ * @handle:	a handle
+ * @fd:		a file descriptor representing that handle
+ *
+ * For ION_IOC_SHARE or ION_IOC_MAP userspace populates the handle field with
+ * the handle returned from ion alloc, and the kernel returns the file
+ * descriptor to share or map in the fd field.  For ION_IOC_IMPORT, userspace
+ * provides the file descriptor and the kernel returns the handle.
+ */
+struct ion_fd_data {
+	ion_user_handle_t handle;
+	int fd;
+};
+
+/**
+ * struct ion_handle_data - a handle passed to/from the kernel
+ * @handle:	a handle
+ */
+struct ion_handle_data {
+	ion_user_handle_t handle;
+};
+
+/**
+ * struct ion_custom_data - metadata passed to/from userspace for a custom ioctl
+ * @cmd:	the custom ioctl function to call
+ * @arg:	additional data to pass to the custom ioctl, typically a user
+ *		pointer to a predefined structure
+ *
+ * This works just like the regular cmd and arg fields of an ioctl.
+ */
+struct ion_custom_data {
+	unsigned int cmd;
+	unsigned long arg;
+};
+
+#define ION_IOC_MAGIC		'I'
+
+/**
+ * DOC: ION_IOC_ALLOC - allocate memory
+ *
+ * Takes an ion_allocation_data struct and returns it with the handle field
+ * populated with the opaque handle for the allocation.
+ */
+#define ION_IOC_ALLOC		_IOWR(ION_IOC_MAGIC, 0, \
+				      struct ion_allocation_data)
+
+/**
+ * DOC: ION_IOC_FREE - free memory
+ *
+ * Takes an ion_handle_data struct and frees the handle.
+ */
+#define ION_IOC_FREE		_IOWR(ION_IOC_MAGIC, 1, struct ion_handle_data)
+
+/**
+ * DOC: ION_IOC_MAP - get a file descriptor to mmap
+ *
+ * Takes an ion_fd_data struct with the handle field populated with a valid
+ * opaque handle.  Returns the struct with the fd field set to a file
+ * descriptor open in the current address space.  This file descriptor
+ * can then be used as an argument to mmap.
+ */
+#define ION_IOC_MAP		_IOWR(ION_IOC_MAGIC, 2, struct ion_fd_data)
+
+/**
+ * DOC: ION_IOC_SHARE - creates a file descriptor to use to share an allocation
+ *
+ * Takes an ion_fd_data struct with the handle field populated with a valid
+ * opaque handle.  Returns the struct with the fd field set to a file
+ * descriptor open in the current address space.  This file descriptor
+ * can then be passed to another process.  The corresponding opaque handle can
+ * be retrieved via ION_IOC_IMPORT.
+ */
+#define ION_IOC_SHARE		_IOWR(ION_IOC_MAGIC, 4, struct ion_fd_data)
+
+/**
+ * DOC: ION_IOC_IMPORT - imports a shared file descriptor
+ *
+ * Takes an ion_fd_data struct with the fd field populated with a valid file
+ * descriptor obtained from ION_IOC_SHARE and returns the struct with the handle
+ * filed set to the corresponding opaque handle.
+ */
+#define ION_IOC_IMPORT		_IOWR(ION_IOC_MAGIC, 5, struct ion_fd_data)
+
+/**
+ * DOC: ION_IOC_SYNC - syncs a shared file descriptors to memory
+ *
+ * Deprecated in favor of using the dma_buf api's correctly (syncing
+ * will happen automatically when the buffer is mapped to a device).
+ * If necessary should be used after touching a cached buffer from the cpu,
+ * this will make the buffer in memory coherent.
+ */
+#define ION_IOC_SYNC		_IOWR(ION_IOC_MAGIC, 7, struct ion_fd_data)
+
+/**
+ * DOC: ION_IOC_CUSTOM - call architecture specific ion ioctl
+ *
+ * Takes the argument of the architecture specific ioctl to call and
+ * passes appropriate userdata for that ioctl
+ */
+#define ION_IOC_CUSTOM		_IOWR(ION_IOC_MAGIC, 6, struct ion_custom_data)
+
+#endif /* _UAPI_LINUX_ION_H */
@@ -0,0 +1,211 @@
+#ifndef _UAPI_MSM_ION_H
+#define _UAPI_MSM_ION_H
+
+#include "ion.h"
+
+enum msm_ion_heap_types {
+	ION_HEAP_TYPE_MSM_START = ION_HEAP_TYPE_CUSTOM + 1,
+	ION_HEAP_TYPE_SECURE_DMA = ION_HEAP_TYPE_MSM_START,
+	ION_HEAP_TYPE_SYSTEM_SECURE,
+	ION_HEAP_TYPE_HYP_CMA,
+	/*
+	 * if you add a heap type here you should also add it to
+	 * heap_types_info[] in msm_ion.c
+	 */
+};
+
+/**
+ * These are the only ids that should be used for Ion heap ids.
+ * The ids listed are the order in which allocation will be attempted
+ * if specified. Don't swap the order of heap ids unless you know what
+ * you are doing!
+ * Id's are spaced by purpose to allow new Id's to be inserted in-between (for
+ * possible fallbacks)
+ */
+
+enum ion_heap_ids {
+	INVALID_HEAP_ID = -1,
+	ION_CP_MM_HEAP_ID = 8,
+	ION_SECURE_HEAP_ID = 9,
+	ION_SECURE_DISPLAY_HEAP_ID = 10,
+	ION_CP_MFC_HEAP_ID = 12,
+	ION_CP_WB_HEAP_ID = 16, /* 8660 only */
+	ION_CAMERA_HEAP_ID = 20, /* 8660 only */
+	ION_SYSTEM_CONTIG_HEAP_ID = 21,
+	ION_ADSP_HEAP_ID = 22,
+	ION_PIL1_HEAP_ID = 23, /* Currently used for other PIL images */
+	ION_SF_HEAP_ID = 24,
+	ION_SYSTEM_HEAP_ID = 25,
+	ION_PIL2_HEAP_ID = 26, /* Currently used for modem firmware images */
+	ION_QSECOM_HEAP_ID = 27,
+	ION_AUDIO_HEAP_ID = 28,
+
+	ION_MM_FIRMWARE_HEAP_ID = 29,
+
+	ION_HEAP_ID_RESERVED = 31 /** Bit reserved for ION_FLAG_SECURE flag */
+};
+
+/*
+ * The IOMMU heap is deprecated! Here are some aliases for backwards
+ * compatibility:
+ */
+#define ION_IOMMU_HEAP_ID ION_SYSTEM_HEAP_ID
+#define ION_HEAP_TYPE_IOMMU ION_HEAP_TYPE_SYSTEM
+
+enum ion_fixed_position {
+	NOT_FIXED,
+	FIXED_LOW,
+	FIXED_MIDDLE,
+	FIXED_HIGH,
+};
+
+enum cp_mem_usage {
+	VIDEO_BITSTREAM = 0x1,
+	VIDEO_PIXEL = 0x2,
+	VIDEO_NONPIXEL = 0x3,
+	DISPLAY_SECURE_CP_USAGE = 0x4,
+	CAMERA_SECURE_CP_USAGE = 0x5,
+	MAX_USAGE = 0x6,
+	UNKNOWN = 0x7FFFFFFF,
+};
+
+/**
+ * Flags to be used when allocating from the secure heap for
+ * content protection
+ */
+#define ION_FLAG_CP_TOUCH (1 << 17)
+#define ION_FLAG_CP_BITSTREAM (1 << 18)
+#define ION_FLAG_CP_PIXEL  (1 << 19)
+#define ION_FLAG_CP_NON_PIXEL (1 << 20)
+#define ION_FLAG_CP_CAMERA (1 << 21)
+#define ION_FLAG_CP_HLOS (1 << 22)
+#define ION_FLAG_CP_HLOS_FREE (1 << 23)
+#define ION_FLAG_CP_SEC_DISPLAY (1 << 25)
+#define ION_FLAG_CP_APP (1 << 26)
+
+/**
+ * Flag to allow non continguous allocation of memory from secure
+ * heap
+ */
+#define ION_FLAG_ALLOW_NON_CONTIG (1 << 24)
+
+/**
+ * Flag to use when allocating to indicate that a heap is secure.
+ */
+#define ION_FLAG_SECURE (1 << ION_HEAP_ID_RESERVED)
+
+/**
+ * Flag for clients to force contiguous memort allocation
+ *
+ * Use of this flag is carefully monitored!
+ */
+#define ION_FLAG_FORCE_CONTIGUOUS (1 << 30)
+
+/*
+ * Used in conjunction with heap which pool memory to force an allocation
+ * to come from the page allocator directly instead of from the pool allocation
+ */
+#define ION_FLAG_POOL_FORCE_ALLOC (1 << 16)
+
+
+#define ION_FLAG_POOL_PREFETCH (1 << 27)
+
+/**
+* Deprecated! Please use the corresponding ION_FLAG_*
+*/
+#define ION_SECURE ION_FLAG_SECURE
+#define ION_FORCE_CONTIGUOUS ION_FLAG_FORCE_CONTIGUOUS
+
+/**
+ * Macro should be used with ion_heap_ids defined above.
+ */
+#define ION_HEAP(bit) (1 << (bit))
+
+#define ION_ADSP_HEAP_NAME	"adsp"
+#define ION_SYSTEM_HEAP_NAME	"system"
+#define ION_VMALLOC_HEAP_NAME	ION_SYSTEM_HEAP_NAME
+#define ION_KMALLOC_HEAP_NAME	"kmalloc"
+#define ION_AUDIO_HEAP_NAME	"audio"
+#define ION_SF_HEAP_NAME	"sf"
+#define ION_MM_HEAP_NAME	"mm"
+#define ION_CAMERA_HEAP_NAME	"camera_preview"
+#define ION_IOMMU_HEAP_NAME	"iommu"
+#define ION_MFC_HEAP_NAME	"mfc"
+#define ION_WB_HEAP_NAME	"wb"
+#define ION_MM_FIRMWARE_HEAP_NAME	"mm_fw"
+#define ION_PIL1_HEAP_NAME  "pil_1"
+#define ION_PIL2_HEAP_NAME  "pil_2"
+#define ION_QSECOM_HEAP_NAME	"qsecom"
+#define ION_SECURE_HEAP_NAME	"secure_heap"
+#define ION_SECURE_DISPLAY_HEAP_NAME "secure_display"
+
+#define ION_SET_CACHED(__cache)		(__cache | ION_FLAG_CACHED)
+#define ION_SET_UNCACHED(__cache)	(__cache & ~ION_FLAG_CACHED)
+
+#define ION_IS_CACHED(__flags)	((__flags) & ION_FLAG_CACHED)
+
+/* struct ion_flush_data - data passed to ion for flushing caches
+ *
+ * @handle:	handle with data to flush
+ * @fd:		fd to flush
+ * @vaddr:	userspace virtual address mapped with mmap
+ * @offset:	offset into the handle to flush
+ * @length:	length of handle to flush
+ *
+ * Performs cache operations on the handle. If p is the start address
+ * of the handle, p + offset through p + offset + length will have
+ * the cache operations performed
+ */
+struct ion_flush_data {
+	ion_user_handle_t handle;
+	int fd;
+	void *vaddr;
+	unsigned int offset;
+	unsigned int length;
+};
+
+struct ion_prefetch_regions {
+	unsigned int vmid;
+	size_t *sizes;
+	unsigned int nr_sizes;
+};
+
+struct ion_prefetch_data {
+	int heap_id;
+	unsigned long len;
+	/* Is unsigned long bad? 32bit compiler vs 64 bit compiler*/
+	struct ion_prefetch_regions *regions;
+	unsigned int nr_regions;
+};
+
+#define ION_IOC_MSM_MAGIC 'M'
+
+/**
+ * DOC: ION_IOC_CLEAN_CACHES - clean the caches
+ *
+ * Clean the caches of the handle specified.
+ */
+#define ION_IOC_CLEAN_CACHES	_IOWR(ION_IOC_MSM_MAGIC, 0, \
+						struct ion_flush_data)
+/**
+ * DOC: ION_IOC_INV_CACHES - invalidate the caches
+ *
+ * Invalidate the caches of the handle specified.
+ */
+#define ION_IOC_INV_CACHES	_IOWR(ION_IOC_MSM_MAGIC, 1, \
+						struct ion_flush_data)
+/**
+ * DOC: ION_IOC_CLEAN_INV_CACHES - clean and invalidate the caches
+ *
+ * Clean and invalidate the caches of the handle specified.
+ */
+#define ION_IOC_CLEAN_INV_CACHES	_IOWR(ION_IOC_MSM_MAGIC, 2, \
+						struct ion_flush_data)
+
+#define ION_IOC_PREFETCH		_IOWR(ION_IOC_MSM_MAGIC, 3, \
+						struct ion_prefetch_data)
+
+#define ION_IOC_DRAIN			_IOWR(ION_IOC_MSM_MAGIC, 4, \
+						struct ion_prefetch_data)
+
+#endif
@@ -0,0 +1,326 @@
+from tinygrad.runtime.ops_dsp import DSPDevice
+
+kernel = """__attribute__((noinline)) void r_6_10_13_4_4_29(float* restrict __attribute__((align_value(128))) data0, const float* restrict __attribute__((align_value(128))) data1, const float* restrict __attribute__((align_value(128))) data2, const float* restrict __attribute__((align_value(128))) data3) {
+  float val0 = data1[0];
+  float val1 = data1[1];
+  float val2 = data1[2];
+  float val3 = data1[3];
+  float val4 = data1[4];
+  float val5 = data1[5];
+  float val6 = data1[6];
+  float val7 = data1[7];
+  float val8 = data1[8];
+  float val9 = data1[9];
+  float val10 = data1[10];
+  float val11 = data1[11];
+  float val12 = data1[12];
+  float val13 = data1[13];
+  float val14 = data1[14];
+  float val15 = data1[15];
+  float val16 = data1[16];
+  float val17 = data1[17];
+  float val18 = data1[18];
+  float val19 = data1[19];
+  float val20 = data1[20];
+  float val21 = data1[21];
+  float val22 = data1[22];
+  float val23 = data1[23];
+  float val24 = data1[24];
+  float val25 = data1[25];
+  float val26 = data1[26];
+  float val27 = data1[27];
+  float val28 = data1[28];
+  for (int ridx0 = 0; ridx0 < 6; ridx0++) {
+    for (int ridx1 = 0; ridx1 < 10; ridx1++) {
+      int alu0 = ((ridx0*1160)+(ridx1*4));
+      float val29 = data3[alu0+1];
+      float val30 = data3[alu0+2];
+      float val31 = data3[alu0+3];
+      float val32 = data3[alu0+40];
+      float val33 = data3[alu0+41];
+      float val34 = data3[alu0+42];
+      float val35 = data3[alu0+43];
+      float val36 = data3[alu0+80];
+      float val37 = data3[alu0+81];
+      float val38 = data3[alu0+82];
+      float val39 = data3[alu0+83];
+      float val40 = data3[alu0+120];
+      float val41 = data3[alu0+121];
+      float val42 = data3[alu0+122];
+      float val43 = data3[alu0+123];
+      float val44 = data3[alu0+160];
+      float val45 = data3[alu0+161];
+      float val46 = data3[alu0+162];
+      float val47 = data3[alu0+163];
+      float val48 = data3[alu0+200];
+      float val49 = data3[alu0+201];
+      float val50 = data3[alu0+202];
+      float val51 = data3[alu0+203];
+      float val52 = data3[alu0+240];
+      float val53 = data3[alu0+241];
+      float val54 = data3[alu0+242];
+      float val55 = data3[alu0+243];
+      float val56 = data3[alu0+280];
+      float val57 = data3[alu0+281];
+      float val58 = data3[alu0+282];
+      float val59 = data3[alu0+283];
+      float val60 = data3[alu0+320];
+      float val61 = data3[alu0+321];
+      float val62 = data3[alu0+322];
+      float val63 = data3[alu0+323];
+      float val64 = data3[alu0+360];
+      float val65 = data3[alu0+361];
+      float val66 = data3[alu0+362];
+      float val67 = data3[alu0+363];
+      float val68 = data3[alu0+400];
+      float val69 = data3[alu0+401];
+      float val70 = data3[alu0+402];
+      float val71 = data3[alu0+403];
+      float val72 = data3[alu0+440];
+      float val73 = data3[alu0+441];
+      float val74 = data3[alu0+442];
+      float val75 = data3[alu0+443];
+      float val76 = data3[alu0+480];
+      float val77 = data3[alu0+481];
+      float val78 = data3[alu0+482];
+      float val79 = data3[alu0+483];
+      float val80 = data3[alu0+520];
+      float val81 = data3[alu0+521];
+      float val82 = data3[alu0+522];
+      float val83 = data3[alu0+523];
+      float val84 = data3[alu0+560];
+      float val85 = data3[alu0+561];
+      float val86 = data3[alu0+562];
+      float val87 = data3[alu0+563];
+      float val88 = data3[alu0+600];
+      float val89 = data3[alu0+601];
+      float val90 = data3[alu0+602];
+      float val91 = data3[alu0+603];
+      float val92 = data3[alu0+640];
+      float val93 = data3[alu0+641];
+      float val94 = data3[alu0+642];
+      float val95 = data3[alu0+643];
+      float val96 = data3[alu0+680];
+      float val97 = data3[alu0+681];
+      float val98 = data3[alu0+682];
+      float val99 = data3[alu0+683];
+      float val100 = data3[alu0+720];
+      float val101 = data3[alu0+721];
+      float val102 = data3[alu0+722];
+      float val103 = data3[alu0+723];
+      float val104 = data3[alu0+760];
+      float val105 = data3[alu0+761];
+      float val106 = data3[alu0+762];
+      float val107 = data3[alu0+763];
+      float val108 = data3[alu0+800];
+      float val109 = data3[alu0+801];
+      float val110 = data3[alu0+802];
+      float val111 = data3[alu0+803];
+      float val112 = data3[alu0+840];
+      float val113 = data3[alu0+841];
+      float val114 = data3[alu0+842];
+      float val115 = data3[alu0+843];
+      float val116 = data3[alu0+880];
+      float val117 = data3[alu0+881];
+      float val118 = data3[alu0+882];
+      float val119 = data3[alu0+883];
+      float val120 = data3[alu0+920];
+      float val121 = data3[alu0+921];
+      float val122 = data3[alu0+922];
+      float val123 = data3[alu0+923];
+      float val124 = data3[alu0+960];
+      float val125 = data3[alu0+961];
+      float val126 = data3[alu0+962];
+      float val127 = data3[alu0+963];
+      float val128 = data3[alu0+1000];
+      float val129 = data3[alu0+1001];
+      float val130 = data3[alu0+1002];
+      float val131 = data3[alu0+1003];
+      float val132 = data3[alu0+1040];
+      float val133 = data3[alu0+1041];
+      float val134 = data3[alu0+1042];
+      float val135 = data3[alu0+1043];
+      float val136 = data3[alu0+1080];
+      float val137 = data3[alu0+1081];
+      float val138 = data3[alu0+1082];
+      float val139 = data3[alu0+1083];
+      float val140 = data3[alu0+1120];
+      float val141 = data3[alu0+1121];
+      float val142 = data3[alu0+1122];
+      float val143 = data3[alu0+1123];
+      float val144 = data3[alu0];
+      for (int ridx2 = 0; ridx2 < 13; ridx2++) {
+        int alu1 = (ridx2*4);
+        int alu2 = ((ridx0*2080)+(ridx1*208)+alu1);
+        float val145 = data2[alu1+1];
+        float cast0 = (float)(((val0!=val145)!=1));
+        float cast1 = (float)(((val1!=val145)!=1));
+        float cast2 = (float)(((val2!=val145)!=1));
+        float cast3 = (float)(((val3!=val145)!=1));
+        float cast4 = (float)(((val4!=val145)!=1));
+        float cast5 = (float)(((val5!=val145)!=1));
+        float cast6 = (float)(((val6!=val145)!=1));
+        float cast7 = (float)(((val7!=val145)!=1));
+        float cast8 = (float)(((val8!=val145)!=1));
+        float cast9 = (float)(((val9!=val145)!=1));
+        float cast10 = (float)(((val10!=val145)!=1));
+        float cast11 = (float)(((val11!=val145)!=1));
+        float cast12 = (float)(((val12!=val145)!=1));
+        float cast13 = (float)(((val13!=val145)!=1));
+        float cast14 = (float)(((val14!=val145)!=1));
+        float cast15 = (float)(((val15!=val145)!=1));
+        float cast16 = (float)(((val16!=val145)!=1));
+        float cast17 = (float)(((val17!=val145)!=1));
+        float cast18 = (float)(((val18!=val145)!=1));
+        float cast19 = (float)(((val19!=val145)!=1));
+        float cast20 = (float)(((val20!=val145)!=1));
+        float cast21 = (float)(((val21!=val145)!=1));
+        float cast22 = (float)(((val22!=val145)!=1));
+        float cast23 = (float)(((val23!=val145)!=1));
+        float cast24 = (float)(((val24!=val145)!=1));
+        float cast25 = (float)(((val25!=val145)!=1));
+        float cast26 = (float)(((val26!=val145)!=1));
+        float cast27 = (float)(((val27!=val145)!=1));
+        float cast28 = (float)(((val28!=val145)!=1));
+        data0[alu2+1] = ((cast0*val144)+(cast1*val32)+(cast2*val36)+(cast3*val40)+(cast4*val44)+(cast5*val48)+(cast6*val52)+(cast7*val56)+(cast8*val60)+(cast9*val64)+(cast10*val68)+(cast11*val72)+(cast12*val76)+(cast13*val80)+(cast14*val84)+(cast15*val88)+(cast16*val92)+(cast17*val96)+(cast18*val100)+(cast19*val104)+(cast20*val108)+(cast21*val112)+(cast22*val116)+(cast23*val120)+(cast24*val124)+(cast25*val128)+(cast26*val132)+(cast27*val136)+(cast28*val140));
+        data0[alu2+53] = ((cast0*val29)+(cast1*val33)+(cast2*val37)+(cast3*val41)+(cast4*val45)+(cast5*val49)+(cast6*val53)+(cast7*val57)+(cast8*val61)+(cast9*val65)+(cast10*val69)+(cast11*val73)+(cast12*val77)+(cast13*val81)+(cast14*val85)+(cast15*val89)+(cast16*val93)+(cast17*val97)+(cast18*val101)+(cast19*val105)+(cast20*val109)+(cast21*val113)+(cast22*val117)+(cast23*val121)+(cast24*val125)+(cast25*val129)+(cast26*val133)+(cast27*val137)+(cast28*val141));
+        data0[alu2+105] = ((cast0*val30)+(cast1*val34)+(cast2*val38)+(cast3*val42)+(cast4*val46)+(cast5*val50)+(cast6*val54)+(cast7*val58)+(cast8*val62)+(cast9*val66)+(cast10*val70)+(cast11*val74)+(cast12*val78)+(cast13*val82)+(cast14*val86)+(cast15*val90)+(cast16*val94)+(cast17*val98)+(cast18*val102)+(cast19*val106)+(cast20*val110)+(cast21*val114)+(cast22*val118)+(cast23*val122)+(cast24*val126)+(cast25*val130)+(cast26*val134)+(cast27*val138)+(cast28*val142));
+        data0[alu2+157] = ((cast0*val31)+(cast1*val35)+(cast2*val39)+(cast3*val43)+(cast4*val47)+(cast5*val51)+(cast6*val55)+(cast7*val59)+(cast8*val63)+(cast9*val67)+(cast10*val71)+(cast11*val75)+(cast12*val79)+(cast13*val83)+(cast14*val87)+(cast15*val91)+(cast16*val95)+(cast17*val99)+(cast18*val103)+(cast19*val107)+(cast20*val111)+(cast21*val115)+(cast22*val119)+(cast23*val123)+(cast24*val127)+(cast25*val131)+(cast26*val135)+(cast27*val139)+(cast28*val143));
+        float val146 = data2[alu1+2];
+        float cast29 = (float)(((val0!=val146)!=1));
+        float cast30 = (float)(((val1!=val146)!=1));
+        float cast31 = (float)(((val2!=val146)!=1));
+        float cast32 = (float)(((val3!=val146)!=1));
+        float cast33 = (float)(((val4!=val146)!=1));
+        float cast34 = (float)(((val5!=val146)!=1));
+        float cast35 = (float)(((val6!=val146)!=1));
+        float cast36 = (float)(((val7!=val146)!=1));
+        float cast37 = (float)(((val8!=val146)!=1));
+        float cast38 = (float)(((val9!=val146)!=1));
+        float cast39 = (float)(((val10!=val146)!=1));
+        float cast40 = (float)(((val11!=val146)!=1));
+        float cast41 = (float)(((val12!=val146)!=1));
+        float cast42 = (float)(((val13!=val146)!=1));
+        float cast43 = (float)(((val14!=val146)!=1));
+        float cast44 = (float)(((val15!=val146)!=1));
+        float cast45 = (float)(((val16!=val146)!=1));
+        float cast46 = (float)(((val17!=val146)!=1));
+        float cast47 = (float)(((val18!=val146)!=1));
+        float cast48 = (float)(((val19!=val146)!=1));
+        float cast49 = (float)(((val20!=val146)!=1));
+        float cast50 = (float)(((val21!=val146)!=1));
+        float cast51 = (float)(((val22!=val146)!=1));
+        float cast52 = (float)(((val23!=val146)!=1));
+        float cast53 = (float)(((val24!=val146)!=1));
+        float cast54 = (float)(((val25!=val146)!=1));
+        float cast55 = (float)(((val26!=val146)!=1));
+        float cast56 = (float)(((val27!=val146)!=1));
+        float cast57 = (float)(((val28!=val146)!=1));
+        data0[alu2+2] = ((cast29*val144)+(cast30*val32)+(cast31*val36)+(cast32*val40)+(cast33*val44)+(cast34*val48)+(cast35*val52)+(cast36*val56)+(cast37*val60)+(cast38*val64)+(cast39*val68)+(cast40*val72)+(cast41*val76)+(cast42*val80)+(cast43*val84)+(cast44*val88)+(cast45*val92)+(cast46*val96)+(cast47*val100)+(cast48*val104)+(cast49*val108)+(cast50*val112)+(cast51*val116)+(cast52*val120)+(cast53*val124)+(cast54*val128)+(cast55*val132)+(cast56*val136)+(cast57*val140));
+        data0[alu2+54] = ((cast29*val29)+(cast30*val33)+(cast31*val37)+(cast32*val41)+(cast33*val45)+(cast34*val49)+(cast35*val53)+(cast36*val57)+(cast37*val61)+(cast38*val65)+(cast39*val69)+(cast40*val73)+(cast41*val77)+(cast42*val81)+(cast43*val85)+(cast44*val89)+(cast45*val93)+(cast46*val97)+(cast47*val101)+(cast48*val105)+(cast49*val109)+(cast50*val113)+(cast51*val117)+(cast52*val121)+(cast53*val125)+(cast54*val129)+(cast55*val133)+(cast56*val137)+(cast57*val141));
+        data0[alu2+106] = ((cast29*val30)+(cast30*val34)+(cast31*val38)+(cast32*val42)+(cast33*val46)+(cast34*val50)+(cast35*val54)+(cast36*val58)+(cast37*val62)+(cast38*val66)+(cast39*val70)+(cast40*val74)+(cast41*val78)+(cast42*val82)+(cast43*val86)+(cast44*val90)+(cast45*val94)+(cast46*val98)+(cast47*val102)+(cast48*val106)+(cast49*val110)+(cast50*val114)+(cast51*val118)+(cast52*val122)+(cast53*val126)+(cast54*val130)+(cast55*val134)+(cast56*val138)+(cast57*val142));
+        data0[alu2+158] = ((cast29*val31)+(cast30*val35)+(cast31*val39)+(cast32*val43)+(cast33*val47)+(cast34*val51)+(cast35*val55)+(cast36*val59)+(cast37*val63)+(cast38*val67)+(cast39*val71)+(cast40*val75)+(cast41*val79)+(cast42*val83)+(cast43*val87)+(cast44*val91)+(cast45*val95)+(cast46*val99)+(cast47*val103)+(cast48*val107)+(cast49*val111)+(cast50*val115)+(cast51*val119)+(cast52*val123)+(cast53*val127)+(cast54*val131)+(cast55*val135)+(cast56*val139)+(cast57*val143));
+        float val147 = data2[alu1+3];
+        float cast58 = (float)(((val0!=val147)!=1));
+        float cast59 = (float)(((val1!=val147)!=1));
+        float cast60 = (float)(((val2!=val147)!=1));
+        float cast61 = (float)(((val3!=val147)!=1));
+        float cast62 = (float)(((val4!=val147)!=1));
+        float cast63 = (float)(((val5!=val147)!=1));
+        float cast64 = (float)(((val6!=val147)!=1));
+        float cast65 = (float)(((val7!=val147)!=1));
+        float cast66 = (float)(((val8!=val147)!=1));
+        float cast67 = (float)(((val9!=val147)!=1));
+        float cast68 = (float)(((val10!=val147)!=1));
+        float cast69 = (float)(((val11!=val147)!=1));
+        float cast70 = (float)(((val12!=val147)!=1));
+        float cast71 = (float)(((val13!=val147)!=1));
+        float cast72 = (float)(((val14!=val147)!=1));
+        float cast73 = (float)(((val15!=val147)!=1));
+        float cast74 = (float)(((val16!=val147)!=1));
+        float cast75 = (float)(((val17!=val147)!=1));
+        float cast76 = (float)(((val18!=val147)!=1));
+        float cast77 = (float)(((val19!=val147)!=1));
+        float cast78 = (float)(((val20!=val147)!=1));
+        float cast79 = (float)(((val21!=val147)!=1));
+        float cast80 = (float)(((val22!=val147)!=1));
+        float cast81 = (float)(((val23!=val147)!=1));
+        float cast82 = (float)(((val24!=val147)!=1));
+        float cast83 = (float)(((val25!=val147)!=1));
+        float cast84 = (float)(((val26!=val147)!=1));
+        float cast85 = (float)(((val27!=val147)!=1));
+        float cast86 = (float)(((val28!=val147)!=1));
+        data0[alu2+3] = ((cast58*val144)+(cast59*val32)+(cast60*val36)+(cast61*val40)+(cast62*val44)+(cast63*val48)+(cast64*val52)+(cast65*val56)+(cast66*val60)+(cast67*val64)+(cast68*val68)+(cast69*val72)+(cast70*val76)+(cast71*val80)+(cast72*val84)+(cast73*val88)+(cast74*val92)+(cast75*val96)+(cast76*val100)+(cast77*val104)+(cast78*val108)+(cast79*val112)+(cast80*val116)+(cast81*val120)+(cast82*val124)+(cast83*val128)+(cast84*val132)+(cast85*val136)+(cast86*val140));
+        data0[alu2+55] = ((cast58*val29)+(cast59*val33)+(cast60*val37)+(cast61*val41)+(cast62*val45)+(cast63*val49)+(cast64*val53)+(cast65*val57)+(cast66*val61)+(cast67*val65)+(cast68*val69)+(cast69*val73)+(cast70*val77)+(cast71*val81)+(cast72*val85)+(cast73*val89)+(cast74*val93)+(cast75*val97)+(cast76*val101)+(cast77*val105)+(cast78*val109)+(cast79*val113)+(cast80*val117)+(cast81*val121)+(cast82*val125)+(cast83*val129)+(cast84*val133)+(cast85*val137)+(cast86*val141));
+        data0[alu2+107] = ((cast58*val30)+(cast59*val34)+(cast60*val38)+(cast61*val42)+(cast62*val46)+(cast63*val50)+(cast64*val54)+(cast65*val58)+(cast66*val62)+(cast67*val66)+(cast68*val70)+(cast69*val74)+(cast70*val78)+(cast71*val82)+(cast72*val86)+(cast73*val90)+(cast74*val94)+(cast75*val98)+(cast76*val102)+(cast77*val106)+(cast78*val110)+(cast79*val114)+(cast80*val118)+(cast81*val122)+(cast82*val126)+(cast83*val130)+(cast84*val134)+(cast85*val138)+(cast86*val142));
+        data0[alu2+159] = ((cast58*val31)+(cast59*val35)+(cast60*val39)+(cast61*val43)+(cast62*val47)+(cast63*val51)+(cast64*val55)+(cast65*val59)+(cast66*val63)+(cast67*val67)+(cast68*val71)+(cast69*val75)+(cast70*val79)+(cast71*val83)+(cast72*val87)+(cast73*val91)+(cast74*val95)+(cast75*val99)+(cast76*val103)+(cast77*val107)+(cast78*val111)+(cast79*val115)+(cast80*val119)+(cast81*val123)+(cast82*val127)+(cast83*val131)+(cast84*val135)+(cast85*val139)+(cast86*val143));
+        float val148 = data2[alu1];
+        float cast87 = (float)(((val0!=val148)!=1));
+        float cast88 = (float)(((val1!=val148)!=1));
+        float cast89 = (float)(((val2!=val148)!=1));
+        float cast90 = (float)(((val3!=val148)!=1));
+        float cast91 = (float)(((val4!=val148)!=1));
+        float cast92 = (float)(((val5!=val148)!=1));
+        float cast93 = (float)(((val6!=val148)!=1));
+        float cast94 = (float)(((val7!=val148)!=1));
+        float cast95 = (float)(((val8!=val148)!=1));
+        float cast96 = (float)(((val9!=val148)!=1));
+        float cast97 = (float)(((val10!=val148)!=1));
+        float cast98 = (float)(((val11!=val148)!=1));
+        float cast99 = (float)(((val12!=val148)!=1));
+        float cast100 = (float)(((val13!=val148)!=1));
+        float cast101 = (float)(((val14!=val148)!=1));
+        float cast102 = (float)(((val15!=val148)!=1));
+        float cast103 = (float)(((val16!=val148)!=1));
+        float cast104 = (float)(((val17!=val148)!=1));
+        float cast105 = (float)(((val18!=val148)!=1));
+        float cast106 = (float)(((val19!=val148)!=1));
+        float cast107 = (float)(((val20!=val148)!=1));
+        float cast108 = (float)(((val21!=val148)!=1));
+        float cast109 = (float)(((val22!=val148)!=1));
+        float cast110 = (float)(((val23!=val148)!=1));
+        float cast111 = (float)(((val24!=val148)!=1));
+        float cast112 = (float)(((val25!=val148)!=1));
+        float cast113 = (float)(((val26!=val148)!=1));
+        float cast114 = (float)(((val27!=val148)!=1));
+        float cast115 = (float)(((val28!=val148)!=1));
+        data0[alu2+52] = ((cast87*val29)+(cast88*val33)+(cast89*val37)+(cast90*val41)+(cast91*val45)+(cast92*val49)+(cast93*val53)+(cast94*val57)+(cast95*val61)+(cast96*val65)+(cast97*val69)+(cast98*val73)+(cast99*val77)+(cast100*val81)+(cast101*val85)+(cast102*val89)+(cast103*val93)+(cast104*val97)+(cast105*val101)+(cast106*val105)+(cast107*val109)+(cast108*val113)+(cast109*val117)+(cast110*val121)+(cast111*val125)+(cast112*val129)+(cast113*val133)+(cast114*val137)+(cast115*val141));
+        data0[alu2+104] = ((cast87*val30)+(cast88*val34)+(cast89*val38)+(cast90*val42)+(cast91*val46)+(cast92*val50)+(cast93*val54)+(cast94*val58)+(cast95*val62)+(cast96*val66)+(cast97*val70)+(cast98*val74)+(cast99*val78)+(cast100*val82)+(cast101*val86)+(cast102*val90)+(cast103*val94)+(cast104*val98)+(cast105*val102)+(cast106*val106)+(cast107*val110)+(cast108*val114)+(cast109*val118)+(cast110*val122)+(cast111*val126)+(cast112*val130)+(cast113*val134)+(cast114*val138)+(cast115*val142));
+        data0[alu2+156] = ((cast87*val31)+(cast88*val35)+(cast89*val39)+(cast90*val43)+(cast91*val47)+(cast92*val51)+(cast93*val55)+(cast94*val59)+(cast95*val63)+(cast96*val67)+(cast97*val71)+(cast98*val75)+(cast99*val79)+(cast100*val83)+(cast101*val87)+(cast102*val91)+(cast103*val95)+(cast104*val99)+(cast105*val103)+(cast106*val107)+(cast107*val111)+(cast108*val115)+(cast109*val119)+(cast110*val123)+(cast111*val127)+(cast112*val131)+(cast113*val135)+(cast114*val139)+(cast115*val143));
+        data0[alu2] = ((cast87*val144)+(cast88*val32)+(cast89*val36)+(cast90*val40)+(cast91*val44)+(cast92*val48)+(cast93*val52)+(cast94*val56)+(cast95*val60)+(cast96*val64)+(cast97*val68)+(cast98*val72)+(cast99*val76)+(cast100*val80)+(cast101*val84)+(cast102*val88)+(cast103*val92)+(cast104*val96)+(cast105*val100)+(cast106*val104)+(cast107*val108)+(cast108*val112)+(cast109*val116)+(cast110*val120)+(cast111*val124)+(cast112*val128)+(cast113*val132)+(cast114*val136)+(cast115*val140));
+      }
+    }
+  }
+}"""
+
+entry = """typedef union { struct { void *pv; unsigned int len; } buf; struct { int fd; unsigned int offset; } dma; } remote_arg;
+void* HAP_mmap(void *addr, int len, int prot, int flags, int fd, long offset);
+int HAP_munmap(void *addr, int len);
+int HAP_mmap_get(int fd, void **vaddr, void **paddr);
+int HAP_mmap_put(int fd);
+unsigned long long HAP_perf_get_time_us(void);
+int entry(unsigned long long handle, unsigned int sc, remote_arg* pra) {
+if ((sc>>24) != 2) return 0;
+
+unsigned long long start = HAP_perf_get_time_us();
+for (int i = 0; i < 50; i++) {
+  void* buf = HAP_mmap(0, 1, 3, 0, pra[2].dma.fd, 0);
+  HAP_munmap(buf, 1);
+}
+*(unsigned long long *)(pra[1].buf.pv) = HAP_perf_get_time_us() - start;
+
+return 0; }
+"""
+
+if __name__ == "__main__":
+  dev = DSPDevice()
+
+  bufs = [dev.allocator.alloc(0x60000) for _ in range(4)]
+
+  only_entry = dev.compiler.compile(entry)
+  app1 = dev.runtime("test", only_entry)
+  x = app1(*bufs)
+
+  entry_n_unsued_code = dev.compiler.compile(kernel + "\n" + entry)
+  app2 = dev.runtime("test", entry_n_unsued_code)
+  x = app2(*bufs)
@@ -0,0 +1,279 @@
+from tinygrad.runtime.ops_dsp import DSPDevice
+
+kernel = """__attribute__((noinline)) void r_64_4_4_64_4_4_4(float* restrict __attribute__((align_value(128))) data0, const float* restrict __attribute__((align_value(128))) data1, const float* restrict __attribute__((align_value(128))) data2, const float* restrict __attribute__((align_value(128))) data3) {
+  for (int ridx0 = 0; ridx0 < 64; ridx0++) {
+    int alu0 = (ridx0*4096);
+    for (int ridx1 = 0; ridx1 < 4; ridx1++) {
+      int alu1 = (ridx1*64);
+      for (int ridx2 = 0; ridx2 < 4; ridx2++) {
+        int alu2 = (ridx2*4);
+        int alu3 = ((ridx0*1024)+alu1+alu2);
+        int alu4 = (alu1+alu2);
+        float val0 = data3[alu4+1];
+        float val1 = data3[alu4+2];
+        float val2 = data3[alu4+3];
+        float val3 = data3[alu4+16];
+        float val4 = data3[alu4+17];
+        float val5 = data3[alu4+18];
+        float val6 = data3[alu4+19];
+        float val7 = data3[alu4+32];
+        float val8 = data3[alu4+33];
+        float val9 = data3[alu4+34];
+        float val10 = data3[alu4+35];
+        float val11 = data3[alu4+48];
+        float val12 = data3[alu4+49];
+        float val13 = data3[alu4+50];
+        float val14 = data3[alu4+51];
+        float val15 = data3[alu4];
+        float acc0 = 0.0f;
+        float acc1 = 0.0f;
+        float acc2 = 0.0f;
+        float acc3 = 0.0f;
+        float acc4 = 0.0f;
+        float acc5 = 0.0f;
+        float acc6 = 0.0f;
+        float acc7 = 0.0f;
+        float acc8 = 0.0f;
+        float acc9 = 0.0f;
+        float acc10 = 0.0f;
+        float acc11 = 0.0f;
+        float acc12 = 0.0f;
+        float acc13 = 0.0f;
+        float acc14 = 0.0f;
+        float acc15 = 0.0f;
+        float acc16 = 0.0f;
+        float acc17 = 0.0f;
+        float acc18 = 0.0f;
+        float acc19 = 0.0f;
+        float acc20 = 0.0f;
+        float acc21 = 0.0f;
+        float acc22 = 0.0f;
+        float acc23 = 0.0f;
+        float acc24 = 0.0f;
+        float acc25 = 0.0f;
+        float acc26 = 0.0f;
+        float acc27 = 0.0f;
+        float acc28 = 0.0f;
+        float acc29 = 0.0f;
+        float acc30 = 0.0f;
+        float acc31 = 0.0f;
+        float acc32 = 0.0f;
+        float acc33 = 0.0f;
+        float acc34 = 0.0f;
+        float acc35 = 0.0f;
+        float acc36 = 0.0f;
+        float acc37 = 0.0f;
+        float acc38 = 0.0f;
+        float acc39 = 0.0f;
+        float acc40 = 0.0f;
+        float acc41 = 0.0f;
+        float acc42 = 0.0f;
+        float acc43 = 0.0f;
+        float acc44 = 0.0f;
+        float acc45 = 0.0f;
+        float acc46 = 0.0f;
+        float acc47 = 0.0f;
+        float acc48 = 0.0f;
+        float acc49 = 0.0f;
+        float acc50 = 0.0f;
+        float acc51 = 0.0f;
+        float acc52 = 0.0f;
+        float acc53 = 0.0f;
+        float acc54 = 0.0f;
+        float acc55 = 0.0f;
+        float acc56 = 0.0f;
+        float acc57 = 0.0f;
+        float acc58 = 0.0f;
+        float acc59 = 0.0f;
+        float acc60 = 0.0f;
+        float acc61 = 0.0f;
+        float acc62 = 0.0f;
+        float acc63 = 0.0f;
+        for (int ridx3 = 0; ridx3 < 64; ridx3++) {
+          int alu5 = (alu0+(ridx2*256)+ridx3);
+          float val16 = data2[alu5+64];
+          float val17 = data2[alu5+128];
+          float val18 = data2[alu5+192];
+          float val19 = data2[alu5+1024];
+          float val20 = data2[alu5+1088];
+          float val21 = data2[alu5+1152];
+          float val22 = data2[alu5+1216];
+          float val23 = data2[alu5+2048];
+          float val24 = data2[alu5+2112];
+          float val25 = data2[alu5+2176];
+          float val26 = data2[alu5+2240];
+          float val27 = data2[alu5+3072];
+          float val28 = data2[alu5+3136];
+          float val29 = data2[alu5+3200];
+          float val30 = data2[alu5+3264];
+          float val31 = data2[alu5];
+          int alu6 = (alu0+(ridx1*256)+ridx3);
+          float val32 = data1[alu6+64];
+          float val33 = data1[alu6+128];
+          float val34 = data1[alu6+192];
+          float val35 = data1[alu6+1024];
+          float val36 = data1[alu6+1088];
+          float val37 = data1[alu6+1152];
+          float val38 = data1[alu6+1216];
+          float val39 = data1[alu6+2048];
+          float val40 = data1[alu6+2112];
+          float val41 = data1[alu6+2176];
+          float val42 = data1[alu6+2240];
+          float val43 = data1[alu6+3072];
+          float val44 = data1[alu6+3136];
+          float val45 = data1[alu6+3200];
+          float val46 = data1[alu6+3264];
+          float val47 = data1[alu6];
+          acc0 = (acc0+(val47*val31));
+          acc1 = (acc1+(val35*val19));
+          acc2 = (acc2+(val39*val23));
+          acc3 = (acc3+(val43*val27));
+          acc4 = (acc4+(val32*val31));
+          acc5 = (acc5+(val36*val19));
+          acc6 = (acc6+(val40*val23));
+          acc7 = (acc7+(val44*val27));
+          acc8 = (acc8+(val33*val31));
+          acc9 = (acc9+(val37*val19));
+          acc10 = (acc10+(val41*val23));
+          acc11 = (acc11+(val45*val27));
+          acc12 = (acc12+(val34*val31));
+          acc13 = (acc13+(val38*val19));
+          acc14 = (acc14+(val42*val23));
+          acc15 = (acc15+(val46*val27));
+          acc16 = (acc16+(val47*val16));
+          acc17 = (acc17+(val35*val20));
+          acc18 = (acc18+(val39*val24));
+          acc19 = (acc19+(val43*val28));
+          acc20 = (acc20+(val32*val16));
+          acc21 = (acc21+(val36*val20));
+          acc22 = (acc22+(val40*val24));
+          acc23 = (acc23+(val44*val28));
+          acc24 = (acc24+(val33*val16));
+          acc25 = (acc25+(val37*val20));
+          acc26 = (acc26+(val41*val24));
+          acc27 = (acc27+(val45*val28));
+          acc28 = (acc28+(val34*val16));
+          acc29 = (acc29+(val38*val20));
+          acc30 = (acc30+(val42*val24));
+          acc31 = (acc31+(val46*val28));
+          acc32 = (acc32+(val47*val17));
+          acc33 = (acc33+(val35*val21));
+          acc34 = (acc34+(val39*val25));
+          acc35 = (acc35+(val43*val29));
+          acc36 = (acc36+(val32*val17));
+          acc37 = (acc37+(val36*val21));
+          acc38 = (acc38+(val40*val25));
+          acc39 = (acc39+(val44*val29));
+          acc40 = (acc40+(val33*val17));
+          acc41 = (acc41+(val37*val21));
+          acc42 = (acc42+(val41*val25));
+          acc43 = (acc43+(val45*val29));
+          acc44 = (acc44+(val34*val17));
+          acc45 = (acc45+(val38*val21));
+          acc46 = (acc46+(val42*val25));
+          acc47 = (acc47+(val46*val29));
+          acc48 = (acc48+(val47*val18));
+          acc49 = (acc49+(val35*val22));
+          acc50 = (acc50+(val39*val26));
+          acc51 = (acc51+(val43*val30));
+          acc52 = (acc52+(val32*val18));
+          acc53 = (acc53+(val36*val22));
+          acc54 = (acc54+(val40*val26));
+          acc55 = (acc55+(val44*val30));
+          acc56 = (acc56+(val33*val18));
+          acc57 = (acc57+(val37*val22));
+          acc58 = (acc58+(val41*val26));
+          acc59 = (acc59+(val45*val30));
+          acc60 = (acc60+(val34*val18));
+          acc61 = (acc61+(val38*val22));
+          acc62 = (acc62+(val42*val26));
+          acc63 = (acc63+(val46*val30));
+        }
+        data0[alu3] = ((acc0*0.125f)+val15);
+        data0[alu3+256] = ((acc1*0.125f)+val15);
+        data0[alu3+512] = ((acc2*0.125f)+val15);
+        data0[alu3+768] = ((acc3*0.125f)+val15);
+        data0[alu3+16] = ((acc4*0.125f)+val3);
+        data0[alu3+272] = ((acc5*0.125f)+val3);
+        data0[alu3+528] = ((acc6*0.125f)+val3);
+        data0[alu3+784] = ((acc7*0.125f)+val3);
+        data0[alu3+32] = ((acc8*0.125f)+val7);
+        data0[alu3+288] = ((acc9*0.125f)+val7);
+        data0[alu3+544] = ((acc10*0.125f)+val7);
+        data0[alu3+800] = ((acc11*0.125f)+val7);
+        data0[alu3+48] = ((acc12*0.125f)+val11);
+        data0[alu3+304] = ((acc13*0.125f)+val11);
+        data0[alu3+560] = ((acc14*0.125f)+val11);
+        data0[alu3+816] = ((acc15*0.125f)+val11);
+        data0[alu3+1] = ((acc16*0.125f)+val0);
+        data0[alu3+257] = ((acc17*0.125f)+val0);
+        data0[alu3+513] = ((acc18*0.125f)+val0);
+        data0[alu3+769] = ((acc19*0.125f)+val0);
+        data0[alu3+17] = ((acc20*0.125f)+val4);
+        data0[alu3+273] = ((acc21*0.125f)+val4);
+        data0[alu3+529] = ((acc22*0.125f)+val4);
+        data0[alu3+785] = ((acc23*0.125f)+val4);
+        data0[alu3+33] = ((acc24*0.125f)+val8);
+        data0[alu3+289] = ((acc25*0.125f)+val8);
+        data0[alu3+545] = ((acc26*0.125f)+val8);
+        data0[alu3+801] = ((acc27*0.125f)+val8);
+        data0[alu3+49] = ((acc28*0.125f)+val12);
+        data0[alu3+305] = ((acc29*0.125f)+val12);
+        data0[alu3+561] = ((acc30*0.125f)+val12);
+        data0[alu3+817] = ((acc31*0.125f)+val12);
+        data0[alu3+2] = ((acc32*0.125f)+val1);
+        data0[alu3+258] = ((acc33*0.125f)+val1);
+        data0[alu3+514] = ((acc34*0.125f)+val1);
+        data0[alu3+770] = ((acc35*0.125f)+val1);
+        data0[alu3+18] = ((acc36*0.125f)+val5);
+        data0[alu3+274] = ((acc37*0.125f)+val5);
+        data0[alu3+530] = ((acc38*0.125f)+val5);
+        data0[alu3+786] = ((acc39*0.125f)+val5);
+        data0[alu3+34] = ((acc40*0.125f)+val9);
+        data0[alu3+290] = ((acc41*0.125f)+val9);
+        data0[alu3+546] = ((acc42*0.125f)+val9);
+        data0[alu3+802] = ((acc43*0.125f)+val9);
+        data0[alu3+50] = ((acc44*0.125f)+val13);
+        data0[alu3+306] = ((acc45*0.125f)+val13);
+        data0[alu3+562] = ((acc46*0.125f)+val13);
+        data0[alu3+818] = ((acc47*0.125f)+val13);
+        data0[alu3+3] = ((acc48*0.125f)+val2);
+        data0[alu3+259] = ((acc49*0.125f)+val2);
+        data0[alu3+515] = ((acc50*0.125f)+val2);
+        data0[alu3+771] = ((acc51*0.125f)+val2);
+        data0[alu3+19] = ((acc52*0.125f)+val6);
+        data0[alu3+275] = ((acc53*0.125f)+val6);
+        data0[alu3+531] = ((acc54*0.125f)+val6);
+        data0[alu3+787] = ((acc55*0.125f)+val6);
+        data0[alu3+35] = ((acc56*0.125f)+val10);
+        data0[alu3+291] = ((acc57*0.125f)+val10);
+        data0[alu3+547] = ((acc58*0.125f)+val10);
+        data0[alu3+803] = ((acc59*0.125f)+val10);
+        data0[alu3+51] = ((acc60*0.125f)+val14);
+        data0[alu3+307] = ((acc61*0.125f)+val14);
+        data0[alu3+563] = ((acc62*0.125f)+val14);
+        data0[alu3+819] = ((acc63*0.125f)+val14);
+      }
+    }
+  }
+}
+"""
+
+entry = """unsigned long long HAP_perf_get_time_us(void);
+int entry(unsigned long long handle, unsigned int sc, void* pra) {
+return HAP_perf_get_time_us() == 1 ? 4 : 0;
+}
+"""
+
+if __name__ == "__main__":
+  dev = DSPDevice()
+
+  bufs = [dev.allocator.alloc(0x60000) for _ in range(4)]
+
+  only_entry = dev.compiler.compile(entry)
+  app1 = dev.runtime("test", only_entry)
+  x = app1(*bufs)
+
+  entry_n_unsued_code = dev.compiler.compile(kernel + "\n" + entry)
+  app2 = dev.runtime("test", entry_n_unsued_code)
+  x = app2(*bufs)
@@ -0,0 +1,151 @@
+#!/usr/bin/env python3
+import os, ctypes, ctypes.util, struct, platform, time
+from tinygrad.runtime.autogen import libc, qcom_dsp
+def to_mv(ptr, sz) -> memoryview: return memoryview(ctypes.cast(ptr, ctypes.POINTER(ctypes.c_uint8 * sz)).contents).cast("B")
+from hexdump import hexdump
+
+def get_struct(argp, stype):
+  return ctypes.cast(ctypes.c_void_p(argp), ctypes.POINTER(stype)).contents
+
+def format_struct(s):
+  sdats = []
+  for field in s._fields_:
+    dat = getattr(s, field[0])
+    if isinstance(dat, int): sdats.append(f"{field[0]}:0x{dat:X}")
+    elif hasattr(dat, "_fields_"): sdats.append((field[0], format_struct(dat)))
+    elif field[0] == "PADDING_0": pass
+    else: sdats.append(f"{field[0]}:{dat}")
+  return sdats
+
+@ctypes.CFUNCTYPE(ctypes.c_int, ctypes.c_int, ctypes.c_ulong, ctypes.c_void_p)
+def ioctl(fd, request, argp):
+  fn = os.readlink(f"/proc/self/fd/{fd}")
+  idir, size, itype, nr = (request>>30), (request>>16)&0x3FFF, (request>>8)&0xFF, request&0xFF
+
+  if fn == "/dev/adsprpc-smd":
+    if nr == 1:
+      st = get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_invoke)
+      method = (st.sc>>24) & 0xFF
+      in_args = (st.sc>>16) & 0xFF
+      out_args = (st.sc>>8) & 0xFF
+      if out_args:
+        for arg in range(in_args, in_args+out_args):
+          ctypes.memset(st.pra[arg].buf.pv, 0, st.pra[arg].buf.len)
+
+  # print("enter", libc.gettid())
+  ret = libc.syscall(0x1d, ctypes.c_int(fd), ctypes.c_ulong(request), ctypes.c_void_p(argp))
+  # print("done", libc.gettid())
+  if fn == "/dev/ion":
+    if nr == 0:
+      st = get_struct(argp, qcom_dsp.struct_ion_allocation_data)
+      print(ret, "ION_IOC_ALLOC", format_struct(st))
+    elif nr == 1:
+      st = get_struct(argp, qcom_dsp.struct_ion_handle_data)
+      print(ret, "ION_IOC_FREE", format_struct(st))
+    elif nr == 2:
+      st = get_struct(argp, qcom_dsp.struct_ion_fd_data)
+      print(ret, "ION_IOC_MAP", format_struct(st))
+  elif fn == "/dev/adsprpc-smd":
+    assert chr(itype) == 'R'
+    if nr == 8:
+      st = ctypes.c_uint32.from_address(argp)
+      print(ret, "FASTRPC_IOCTL_GETINFO", st.value)
+    elif nr == 2:
+      st = get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_mmap)
+      print(ret, "FASTRPC_IOCTL_MMAP", format_struct(st))
+    elif nr == 1:
+      # https://research.checkpoint.com/2021/pwn2own-qualcomm-dsp/
+      st = get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_invoke)
+      print(ret, "FASTRPC_IOCTL_INVOKE", format_struct(st))
+      # 0xFF000000 = Method index and attribute (the highest byte)
+      # 0x00FF0000 = Number of input arguments
+      # 0x0000FF00 = Number of output arguments
+      # 0x000000F0 = Number of input handles
+      # 0x0000000F = Number of output handles
+
+      method = (st.sc>>24) & 0xFF
+      in_args = (st.sc>>16) & 0xFF
+      out_args = (st.sc>>8) & 0xFF
+      in_h = (st.sc>>4) & 0xF
+      out_h = (st.sc>>0) & 0xF
+      print(f"\tm:{method} ia:{in_args} oa:{out_args} ih:{in_h} oh:{out_h}")
+      if in_args or out_args:
+        for arg in range(in_args+out_args):
+          print(arg, format_struct(st.pra[arg]))
+          if st.pra[arg].buf.pv is not None:
+            ww = to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)
+            hexdump(to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)[:0x40])
+    elif nr == 6:
+      print(ret, "FASTRPC_IOCTL_INIT", format_struct(ini:=get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_init)))
+      print(os.readlink(f"/proc/self/fd/{ini.filefd}"))
+      # print(bytearray(to_mv(ini.file, ini.filelen)))
+    elif nr == 7:
+      print(ret, "FASTRPC_IOCTL_INVOKE_ATTRS", format_struct(ini:=get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_invoke_attrs)))
+    elif nr == 12: print(ret, "FASTRPC_IOCTL_CONTROL", format_struct(get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_control)))
+    else:
+      print(f"{ret} UNPARSED {nr}")
+  else:
+    print("ioctl", f"{idir=} {size=} {itype=} {nr=} {fd=} {ret=}", fn)
+  return ret
+
+def install_hook(c_function, python_function):
+  orig_func = (ctypes.c_char*4096)()
+  python_function_addr = ctypes.cast(ctypes.byref(python_function), ctypes.POINTER(ctypes.c_ulong)).contents.value
+  # AARCH64 trampoline to ioctl
+  # 0x0000000000000000:  70 00 00 10    adr x16, #0xc
+  # 0x0000000000000004:  10 02 40 F9    ldr x16, [x16]
+  # 0x0000000000000008:  00 02 1F D6    br  x16
+  tramp = b"\x70\x00\x00\x10\x10\x02\x40\xf9\x00\x02\x1f\xd6"
+  tramp += struct.pack("Q", python_function_addr)
+
+  # get real ioctl address
+  ioctl_address = ctypes.cast(ctypes.byref(c_function), ctypes.POINTER(ctypes.c_ulong))
+
+  # hook ioctl
+  ret = libc.mprotect(ctypes.c_ulong((ioctl_address.contents.value//0x1000)*0x1000), 0x2000, 7)
+  assert ret == 0
+  ret = libc.mprotect(ctypes.c_ulong((ctypes.addressof(orig_func)//0x1000)*0x1000), 0x3000, 7)
+  assert ret == 0
+  libc.memcpy(orig_func, ioctl_address.contents, 0x1000)
+  libc.memcpy(ioctl_address.contents, ctypes.create_string_buffer(tramp), len(tramp))
+  return orig_func
+
+libc = ctypes.CDLL(ctypes.util.find_library("libc"))
+install_hook(libc.ioctl, ioctl)
+adsp = ctypes.CDLL(ctypes.util.find_library("adsprpc"))
+
+def send_rpc_invoke(filename):
+  pass
+
+if __name__ == "__main__":
+  print("calculator_open")
+  # /dsp/cdsp/fastrpc_shell_3
+
+  handle = ctypes.c_int64(-1)
+  z = adsp.remote_handle64_open(ctypes.create_string_buffer(b"file:///libcalculator_skel.so?calculator_skel_handle_invoke&_modver=1.0&_dom=cdsp"),
+                            ctypes.byref(handle))
+
+  print("handle", z, hex(handle.value))
+  assert handle.value != -1
+  test = (ctypes.c_int32 * 100)()
+  for i in range(100): test[i] = i
+  print("calculator_sum")
+  pra = (qcom_dsp.union_remote_arg64 * 3)()
+  #arg_0 = ctypes.c_int32(100)
+  arg_0 = ctypes.c_int32(100)
+  arg_2 = ctypes.c_int64(-1)
+  pra[0].buf.pv = ctypes.addressof(arg_0)
+  pra[0].buf.len = 4
+  pra[1].buf.pv = ctypes.addressof(test)
+  pra[1].buf.len = 0x190
+  pra[2].buf.pv = ctypes.addressof(arg_2)
+  pra[2].buf.len = 8
+  adsp.remote_handle64_invoke(handle, (2<<24) | (2<<16) | (1<<8), pra)
+  print(arg_2.value)
+  print("done")
+
+  print("closing")
+  x = adsp.remote_handle64_close(handle)
+  print(x)
+  print("dun")
+  os._exit(0)
@@ -0,0 +1,312 @@
+#!/usr/bin/env python3
+import os, ctypes, ctypes.util, struct, platform, pathlib, contextlib, mmap, array
+from threading import Thread
+from tinygrad.runtime.autogen import qcom_dsp
+from tinygrad.helpers import round_up, mv_address, to_mv
+from hexdump import hexdump
+
+def get_struct(argp, stype):
+  return ctypes.cast(ctypes.c_void_p(argp), ctypes.POINTER(stype)).contents
+
+def format_struct(s):
+  sdats = []
+  for field in s._fields_:
+    dat = getattr(s, field[0])
+    if isinstance(dat, int): sdats.append(f"{field[0]}:0x{dat:X}")
+    elif hasattr(dat, "_fields_"): sdats.append((field[0], format_struct(dat)))
+    elif field[0] == "PADDING_0": pass
+    else: sdats.append(f"{field[0]}:{dat}")
+  return sdats
+
+@ctypes.CFUNCTYPE(ctypes.c_int, ctypes.c_int, ctypes.c_ulong, ctypes.c_void_p)
+def ioctl(fd, request, argp):
+  fn = os.readlink(f"/proc/self/fd/{fd}")
+  idir, size, itype, nr = (request>>30), (request>>16)&0x3FFF, (request>>8)&0xFF, request&0xFF
+
+  # print("enter", libc.gettid())
+  ret = libc.syscall(0x1d, ctypes.c_int(fd), ctypes.c_ulong(request), ctypes.c_void_p(argp))
+  # print("done", libc.gettid())
+  if fn == "/dev/ion":
+    if nr == 0:
+      st = get_struct(argp, qcom_dsp.struct_ion_allocation_data)
+      print(ret, "ION_IOC_ALLOC", format_struct(st))
+    elif nr == 1:
+      st = get_struct(argp, qcom_dsp.struct_ion_handle_data)
+      print(ret, "ION_IOC_FREE", format_struct(st))
+    elif nr == 2:
+      st = get_struct(argp, qcom_dsp.struct_ion_fd_data)
+      print(ret, "ION_IOC_MAP", format_struct(st))
+  elif fn == "/dev/adsprpc-smd":
+    assert chr(itype) == 'R'
+    if nr == 8:
+      st = ctypes.c_uint32.from_address(argp)
+      print(ret, "FASTRPC_IOCTL_GETINFO", st.value)
+    elif nr == 2:
+      st = get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_mmap)
+      print(ret, "FASTRPC_IOCTL_MMAP", format_struct(st))
+    elif nr == 1:
+      # https://research.checkpoint.com/2021/pwn2own-qualcomm-dsp/
+      st = get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_invoke)
+      print(ret, "FASTRPC_IOCTL_INVOKE", format_struct(st))
+      # 0xFF000000 = Method index and attribute (the highest byte)
+      # 0x00FF0000 = Number of input arguments
+      # 0x0000FF00 = Number of output arguments
+      # 0x000000F0 = Number of input handles
+      # 0x0000000F = Number of output handles
+
+      method = (st.sc>>24) & 0xFF
+      in_args = (st.sc>>16) & 0xFF
+      out_args = (st.sc>>8) & 0xFF
+      in_h = (st.sc>>4) & 0xF
+      out_h = (st.sc>>0) & 0xF
+      print(f"\tm:{method} ia:{in_args} oa:{out_args} ih:{in_h} oh:{out_h}")
+      if in_args or out_args:
+        for arg in range(in_args+out_args):
+          print(arg, format_struct(st.pra[arg]))
+          # print(arg, f"arg (0x{st.pra[arg].buf.pv:X} len=0x{st.pra[arg].buf.len:X})")
+          # print("input" if arg < in_args else "output", f"arg (0x{st.pra[arg].buf.pv:X} len=0x{st.pra[arg].buf.len:X})")
+          if st.pra[arg].buf.pv is not None:
+            # if st.pra[arg].buf.len == 0x258:
+            #   print(bytearray(to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)))
+            if st.pra[arg].buf.len == 0x68:
+              print(bytearray(to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)))
+
+            cut = 0x2000 if st.pra[arg].buf.len == 0x2000 or st.pra[arg].buf.len == 0x258 else 0x100
+            ww = to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)
+            hexdump(to_mv(st.pra[arg].buf.pv, st.pra[arg].buf.len)[:cut])
+
+            # if st.pra[arg].buf.len == 0x1000 and ww[0x30] == 0x6e:
+            #   z = ww.cast('Q')[1] + 0x7F00000000
+            #   print("DOO")
+            #   hexdump(to_mv(z, 0x200))
+      #print(format_struct(st.pra)))
+    elif nr == 6:
+      print(ret, "FASTRPC_IOCTL_INIT", format_struct(ini:=get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_init)))
+      print(os.readlink(f"/proc/self/fd/{ini.filefd}"))
+      # print(bytearray(to_mv(ini.file, ini.filelen)))
+    elif nr == 7:
+      print(ret, "FASTRPC_IOCTL_INVOKE_ATTRS", format_struct(ini:=get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_invoke_attrs)))
+    elif nr == 12: print(ret, "FASTRPC_IOCTL_CONTROL", format_struct(get_struct(argp, qcom_dsp.struct_fastrpc_ioctl_control)))
+    else:
+      print(f"{ret} UNPARSED {nr}")
+  else:
+    print("ioctl", f"{idir=} {size=} {itype=} {nr=} {fd=} {ret=}", fn)
+  return ret
+
+def install_hook(c_function, python_function):
+  orig_func = (ctypes.c_char*4096)()
+  python_function_addr = ctypes.cast(ctypes.byref(python_function), ctypes.POINTER(ctypes.c_ulong)).contents.value
+  # AARCH64 trampoline to ioctl
+  # 0x0000000000000000:  70 00 00 10    adr x16, #0xc
+  # 0x0000000000000004:  10 02 40 F9    ldr x16, [x16]
+  # 0x0000000000000008:  00 02 1F D6    br  x16
+  tramp = b"\x70\x00\x00\x10\x10\x02\x40\xf9\x00\x02\x1f\xd6"
+  tramp += struct.pack("Q", python_function_addr)
+
+  # get real ioctl address
+  ioctl_address = ctypes.cast(ctypes.byref(c_function), ctypes.POINTER(ctypes.c_ulong))
+
+  # hook ioctl
+  ret = libc.mprotect(ctypes.c_ulong((ioctl_address.contents.value//0x1000)*0x1000), 0x2000, 7)
+  assert ret == 0
+  ret = libc.mprotect(ctypes.c_ulong((ctypes.addressof(orig_func)//0x1000)*0x1000), 0x3000, 7)
+  assert ret == 0
+  libc.memcpy(orig_func, ioctl_address.contents, 0x1000)
+  libc.memcpy(ioctl_address.contents, ctypes.create_string_buffer(tramp), len(tramp))
+  return orig_func
+
+libc = ctypes.CDLL(ctypes.util.find_library("libc"))
+install_hook(libc.ioctl, ioctl)
+from tinygrad.runtime.autogen import libc
+
+# adsp = ctypes.CDLL(ctypes.util.find_library("adsprpc"))
+# print(adsp)
+
+def rpc_invoke(rpcfd, handle, method, ins=None, outs=None):
+  if ins or outs:
+    ins = ins or list()
+    outs = outs or list()
+    pra = (qcom_dsp.union_remote_arg * (len(ins) + len(outs)))()
+    for i,mv in enumerate(ins + outs):
+      if isinstance(mv, memoryview):
+        pra[i].buf.pv = mv_address(mv) if mv.nbytes > 0 else 0
+        pra[i].buf.len = mv.nbytes
+      else: assert False, "not supported"
+    # pra = (qcom_dsp.union_remote_arg * (len(ins) + len(outs))).from_address(ctypes.addressof(pra))
+  else:
+    pra = None
+    ins = ins or list()
+    outs = outs or list()
+
+  sc = (method << 24) | (len(ins) << 16) | (len(outs) << 8)
+  return qcom_dsp.FASTRPC_IOCTL_INVOKE(rpcfd, handle=handle, sc=sc, pra=pra)
+
+def listner_worker():
+  context = 0
+  handle = 0xffffffff
+  msg_send = memoryview(bytearray(0x10)).cast('I')
+  msg_recv = memoryview(bytearray(0x10)).cast('I')
+  out_buf = memoryview(bytearray(0x1000)).cast('I')
+  in_buf = memoryview(bytearray(0x1000)).cast('I')
+
+  prev_res = 0xffffffff
+  out_buf_size = 0
+
+  req_args = (qcom_dsp.union_remote_arg * 4)()
+  req_args[0].buf = qcom_dsp.struct_remote_buf(pv=mv_address(msg_send), len=0x10)
+  req_args[1].buf = qcom_dsp.struct_remote_buf(pv=mv_address(out_buf), len=0x1000)
+  req_args[2].buf = qcom_dsp.struct_remote_buf(pv=mv_address(msg_recv), len=0x10)
+  req_args[3].buf = qcom_dsp.struct_remote_buf(pv=mv_address(in_buf), len=0x1000)
+
+  while True:
+    msg_send[0] = context
+    msg_send[1] = prev_res
+    msg_send[2] = out_buf_size
+    msg_send[3] = 0x1000
+
+    req_args[1].buf.len = out_buf_size
+    qcom_dsp.FASTRPC_IOCTL_INVOKE(rpcfd, handle=0x3, sc=0x04020200, pra=req_args) # listener
+
+    context = msg_recv[0]
+    handle = msg_recv[1]
+    sc = msg_recv[2]
+    inbufs = (sc >> 16) & 0xff
+    outbufs = (sc >> 8) & 0xff
+
+    in_args, out_args = [], []
+    ptr = mv_address(in_buf)
+    for i in range(inbufs):
+      sz = to_mv(ptr, 4).cast('I')[0]
+      obj_ptr = round_up(ptr + 4, 8)
+      in_args.append(to_mv(obj_ptr, sz))
+      ptr = obj_ptr + sz
+
+    ctypes.memset(mv_address(out_buf), 0, 0x1000)
+    ptr_out = mv_address(out_buf)
+    for i in range(outbufs):
+      sz = to_mv(ptr, 4).cast('I')[0]
+      ptr += 4
+
+      to_mv(ptr_out, 4).cast('I')[0] = sz
+      obj_ptr = round_up(ptr_out + 4, 8)
+
+      out_args.append(to_mv(obj_ptr, sz))
+      ptr_out = obj_ptr + sz
+
+    out_buf_size = ptr_out - mv_address(out_buf)
+
+    if sc == 0x20200: # greating?
+      prev_res = 0
+    elif sc == 0x13050100: # open
+      # for a in in_args: hexdump(a)
+      try:
+        fd = os.open(in_args[3].tobytes()[:-1].decode(), os.O_RDONLY)
+        out_args[0].cast('I')[0] = fd
+        prev_res = 0
+      except: prev_res = 2
+    elif sc == 0x9010000: # seek
+      res = os.lseek(in_args[0].cast('I')[0], in_args[0].cast('I')[1], in_args[0].cast('I')[2])
+      prev_res = 0 if res >= 0 else res
+    elif sc == 0x4010200: # read
+      buf = os.read(in_args[0].cast('I')[0], in_args[0].cast('I')[1])
+      out_args[1][:len(buf)] = buf
+      out_args[0].cast('I')[0] = len(buf)
+      out_args[0].cast('I')[1] = int(len(buf) == 0)
+      prev_res = 0
+    elif sc == 0x3010000: # close
+      os.close(in_args[0].cast('I')[0])
+      prev_res = 0
+    elif sc == 0x1f020100: # stat
+      # try:
+      stat = os.stat(in_args[1].tobytes()[:-1].decode())
+      out_stat = out_args[0].cast('Q')
+      out_stat[1] = stat.st_dev
+      out_stat[2] = stat.st_ino
+      out_stat[3] = stat.st_mode | (stat.st_nlink << 32)
+      out_stat[4] = stat.st_rdev
+      out_stat[5] = stat.st_size
+      # print(stat, stat.st_rdev)
+      # assert False
+      prev_res = 0
+      # except: prev_res = 2
+    elif sc == 0x2010100:
+      heapid = in_args[0].cast('I')[0]
+      lflags = in_args[0].cast('I')[1]
+      rflags = in_args[0].cast('I')[2]
+      assert rflags == 0x1000
+
+      # print(in_args[0])
+
+      # print("WOOW", in_args[0].cast('Q')[2])
+      # print("WOOW2", in_args[0].cast('Q')[2])
+      # print("WOOW3", in_args[0].cast('Q')[3])
+      # print("WOOW3", in_args[0].cast('Q')[3])
+
+      vin = in_args[0].cast('Q')[2]
+      sz = in_args[0].cast('Q')[3]
+      # vin = to_mv(in_args[0].cast('Q')[2], 8).cast('Q')[0]
+      # sz = to_mv(in_args[0].cast('Q')[3], 8).cast('Q')[0]
+
+      st = qcom_dsp.FASTRPC_IOCTL_MMAP(rpcfd, fd=-1, flags=rflags, vaddrin=0, size=sz)
+      out_args[0].cast('Q')[0] = 0
+      out_args[0].cast('Q')[1] = st.vaddrout
+      prev_res = 0
+    else: raise RuntimeError(f"Unknown {sc=:X}")
+
+if __name__ == "__main__":
+  ionfd = os.open('/dev/ion', os.O_RDONLY)
+  rpcfd = os.open('/dev/adsprpc-smd', os.O_RDONLY | os.O_NONBLOCK)
+
+  with contextlib.suppress(RuntimeError, OSError): qcom_dsp.ION_IOC_FREE(ionfd, handle=0)
+  info = qcom_dsp.FASTRPC_IOCTL_GETINFO(rpcfd, 3)
+  # x = qcom_dsp.FASTRPC_IOCTL_SETMODE(rpcfd, 0, __force_as_val=True)
+
+  # init shell?
+  fastrpc_shell = memoryview(bytearray(pathlib.Path('/vendor/dsp/cdsp/fastrpc_shell_3').read_bytes()))
+  shell_mem = qcom_dsp.ION_IOC_ALLOC(ionfd, len=round_up(fastrpc_shell.nbytes, 0x1000), align=0x1000, heap_id_mask=0x2000000, flags=0x1)
+  shell_mapped = qcom_dsp.ION_IOC_MAP(ionfd, handle=shell_mem.handle)
+  fastrpc_shell_addr = libc.mmap(0, shell_mem.len, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED, shell_mapped.fd, 0)
+
+  ctypes.memmove(fastrpc_shell_addr, mv_address(fastrpc_shell), fastrpc_shell.nbytes)
+  # ctypes.memset(fastrpc_shell_addr, 0x0, 0xd6000)
+  # print(hex(fastrpc_shell_addr))
+
+  ctrls = qcom_dsp.FASTRPC_IOCTL_CONTROL(rpcfd, req=0x3)
+
+  init = qcom_dsp.FASTRPC_IOCTL_INIT(rpcfd, flags=0x1, file=fastrpc_shell_addr, filelen=fastrpc_shell.nbytes, filefd=shell_mapped.fd)
+  print("init shell done", shell_mapped.fd)
+
+  # TODO: unmap here
+  # qcom_dsp.ION_IOC_FREE(ionfd, handle=shell_mem.handle)
+
+  rpc_invoke(rpcfd, handle=3, method=3)
+
+  thread = Thread(target=listner_worker)
+  thread.start()
+
+  a1 = memoryview(bytearray(b'\x52\x00\x00\x00\xFF\x00\x00\x00'))
+  a2 = memoryview(bytearray(b"file:///libcalculator_skel.so?calculator_skel_handle_invoke&_modver=1.0&_dom=cdsp\0"))
+  o1 = memoryview(bytearray(0x8))
+  o2 = memoryview(bytearray(0xff))
+  z = rpc_invoke(rpcfd, handle=0, method=0, ins=[a1, a2], outs=[o1, o2])
+  prg_handle = o1.cast('I')[0]
+
+  # test
+  test = (ctypes.c_int32 * 100)()
+  for i in range(100): test[i] = i
+  print("calculator_sum")
+  pra = (qcom_dsp.union_remote_arg * 3)()
+  #arg_0 = ctypes.c_int32(100)
+  arg_0 = ctypes.c_int32(100)
+  arg_2 = ctypes.c_int64(-1)
+  pra[0].buf.pv = ctypes.addressof(arg_0)
+  pra[0].buf.len = 4
+  pra[1].buf.pv = ctypes.addressof(test)
+  pra[1].buf.len = 0x190
+  pra[2].buf.pv = ctypes.addressof(arg_2)
+  pra[2].buf.len = 8
+  qcom_dsp.FASTRPC_IOCTL_INVOKE(rpcfd, handle=prg_handle, sc=(2<<24) | (2<<16) | (1<<8), pra=pra)
+
+  print(arg_2.value)
+  print("done")
+  os._exit(0)
@@ -26,9 +26,9 @@ __kernel void test(__global float* data0, const __global int* data1, const __glo
 """))
 #with open("/tmp/test.elf", "wb") as f: f.write(prog.lib)

-a = Buffer("GPU", 8, dtypes.float32)
-b = Buffer("GPU", 0x10, dtypes.float16)
-c = Buffer("GPU", 8*0x10, dtypes.float16)
+a = Buffer("GPU", 8, dtypes.float32).allocate()
+b = Buffer("GPU", 0x10, dtypes.float16).allocate()
+c = Buffer("GPU", 8*0x10, dtypes.float16).allocate()

 row = np.array([1,2,3,4,5,6,7,8,1,2,3,4,5,6,7,8], np.float16)
 mat = np.random.random((8, 0x10)).astype(np.float16)
@@ -3,6 +3,9 @@ from tinygrad.helpers import getenv
 from tinygrad import dtypes, Tensor
 dtype_in = dtypes.half if getenv("HALF") else dtypes.bfloat16 if getenv("BFLOAT16") else dtypes.float
 acc_dtype = dtypes.half if getenv("ACC_HALF") else dtypes.bfloat16 if getenv("ACC_BFLOAT16") else None
+if getenv("INT"):
+  dtype_in = dtypes.int8
+  acc_dtype = dtypes.int32
 N = getenv("N", 4096)
 M = getenv("M", N)
 K = getenv("K", N)
@@ -87,7 +87,7 @@ if __name__ == "__main__":
  assert '.extern .shared' not in src
  prg = Program("matmul_kernel", src, dname=Device.DEFAULT,
                global_size=[M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1], local_size=[32*compiled.metadata.num_warps, 1, 1],
-                op_estimate=2*M*K*N, mem_estimate=A.nbytes() + B.nbytes() + C.nbytes())
+                mem_estimate=A.nbytes() + B.nbytes() + C.nbytes())
  ei = ExecItem(CompiledRunner(prg), [x.ensure_allocated() for x in si.bufs], si.metadata)
  tflops = []
  for i in range(5):
--- a/Show More
+++ b/Show More