feat: results

* gradient: add TRUNC backward

* test: move round quantization gradient to test_ops

.github/workflows/autogen.yml

@@ -33,12 +33,8 @@ jobs:
       uses: ./.github/actions/setup-tinygrad
       with:
         key: 'autogen'
-        opencl: 'true'
         amd: 'true'
-        cuda: 'true'
         llvm: 'true'
-        webgpu: 'true'
-        mesa: 'true'
         pydeps: 'pyyaml mako'
     - name: Install autogen support packages
       run: sudo apt-get install -y --no-install-recommends libclang-20-dev llvm-20-dev hip-dev libusb-1.0-0-dev libdrm-dev
@@ -48,7 +44,7 @@ jobs:
         python3 -c "from tinygrad.runtime.autogen import opencl"
         python3 -c "from tinygrad.runtime.autogen import cuda, nvrtc, nvjitlink, nv_570, nv_580, nv"
         python3 -c "from tinygrad.runtime.autogen import comgr_3, hsa, hip, amd_gpu, sqtt, rocprof, amdgpu_kd, amdgpu_drm"
-        python3 -c "from tinygrad.runtime.autogen.am import am, pm4_soc15, pm4_nv, sdma_4_0_0, sdma_5_0_0, sdma_6_0_0, smu_v13_0_0, smu_v13_0_6, smu_v13_0_12, smu_v14_0_2"
+        python3 -c "from tinygrad.runtime.autogen.am import *"
         python3 -c "from tinygrad.runtime.autogen import libc, kfd, io_uring, ib, pci, vfio"
         python3 -c "from tinygrad.runtime.autogen import llvm"
         python3 -c "from tinygrad.runtime.autogen import webgpu"
@@ -58,6 +54,7 @@ jobs:
         python3 -c "from tinygrad.runtime.autogen import avcodec"
         python3 -c "from tinygrad.runtime.autogen import llvm_qcom"
         python3 -c "from tinygrad.runtime.autogen import mlx5"
+        python3 -c "from tinygrad.runtime.autogen import ggml_common"
         REGEN=1 python3 -c "from tinygrad.runtime.autogen import libclang"
     - name: Check for differences
       run: |

.github/workflows/benchmark.yml

@@ -51,34 +51,35 @@ jobs:
     - name: openpilot compile3 0.10.1 driving_vision
       run: FLOAT16=1 DEV=CL IMAGE=1 python3.11 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
 
-  testframeworkpytest:
-    name: framework pytest
-    env:
-      CI: ""
-      CAPTURE_PROCESS_REPLAY: "0"
-    runs-on: [self-hosted, framework]
-    timeout-minutes: 10
-    defaults:
-      run:
-        shell: bash -e -o pipefail {0}
-    if: github.repository_owner == 'tinygrad'
-    steps:
-    - name: Checkout Code
-      uses: actions/checkout@v6
-    - name: setup python environment
-      run: |
-        rm -rf /tmp/tinygrad_pytest_ci
-        uv venv /tmp/tinygrad_pytest_ci
-        source /tmp/tinygrad_pytest_ci/bin/activate
-        uv pip install .[testing]
-    - name: setup staging db
-      run: |
-        echo "CACHEDB=/tmp/pytest-db-ci.db" >> $GITHUB_ENV
-        rm -f /tmp/pytest-db-ci*
-    - name: Run pytest -nauto
-      run: |
-        source /tmp/tinygrad_pytest_ci/bin/activate
-        pytest -nauto --durations=20
+  # TODO: reenable when not flaky
+  #testframeworkpytest:
+  #  name: framework pytest
+  #  env:
+  #    CI: ""
+  #    CAPTURE_PROCESS_REPLAY: "0"
+  #  runs-on: [self-hosted, framework]
+  #  timeout-minutes: 10
+  #  defaults:
+  #    run:
+  #      shell: bash -e -o pipefail {0}
+  #  if: github.repository_owner == 'tinygrad'
+  #  steps:
+  #  - name: Checkout Code
+  #    uses: actions/checkout@v6
+  #  - name: setup python environment
+  #    run: |
+  #      rm -rf /tmp/tinygrad_pytest_ci
+  #      uv venv /tmp/tinygrad_pytest_ci
+  #      source /tmp/tinygrad_pytest_ci/bin/activate
+  #      uv pip install .[testing]
+  #  - name: setup staging db
+  #    run: |
+  #      echo "CACHEDB=/tmp/pytest-db-ci.db" >> $GITHUB_ENV
+  #      rm -f /tmp/pytest-db-ci*
+  #  - name: Run pytest -nauto
+  #    run: |
+  #      source /tmp/tinygrad_pytest_ci/bin/activate
+  #      pytest -nauto --durations=20
 
   testmacbenchmark:
     name: Mac Benchmark
@@ -189,7 +190,7 @@ jobs:
         path: |
           onnx_inference_speed.csv
     - 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.11 process_replay.py
+      uses: ./.github/actions/process-replay
 
   testusbgpu:
     name: UsbGPU Benchmark
@@ -324,7 +325,7 @@ jobs:
         path: |
           onnx_inference_speed.csv
     - 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: ./.github/actions/process-replay
 
   testmorenvidiabenchmark:
     name: tinybox green Training Benchmark
@@ -386,7 +387,7 @@ jobs:
       # TODO: remove BERT_LAYERS once scheduler is fast
       run: BENCHMARK_LOG=bert_10steps_6gpu DEV=NV CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=72 GPUS=6 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
     - 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: ./.github/actions/process-replay
 
   testamdbenchmark:
     name: tinybox red Benchmark
@@ -498,7 +499,7 @@ jobs:
     - name: Run GPT2 w HALF/BEAM
       run: BENCHMARK_LOG=gpt2_half_beam DEV=AMD HALF=1 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 examples/gpt2.py --count 10 --temperature 0 --timing
     - 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: ./.github/actions/process-replay
 
   testmoreamdbenchmark:
     name: tinybox red Training Benchmark
@@ -555,7 +556,7 @@ jobs:
     #- name: Test full tinyfs load
     #  run: TINYFS_ENDPOINT=10.0.52.11:6767 PYTHONPATH=. python extra/tinyfs/fetch_file.py --hash d734f5e3be9f1e9d863bfaa4fc6c1ef2 --len 175866113 --dest mapping.json --check
     - 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: ./.github/actions/process-replay
 
   testmlperfamdbenchmark:
     name: tinybox red MLPerf Benchmark
@@ -601,7 +602,7 @@ jobs:
       # TODO: remove BERT_LAYERS once scheduler is fast
       run: BENCHMARK_LOG=bert_10steps_6gpu DEV=AMD CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=72 GPUS=6 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py
     - 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: ./.github/actions/process-replay
 
   testqualcommbenchmark:
     name: comma Benchmark
@@ -623,6 +624,8 @@ jobs:
       run: test/external/process_replay/reset.py
     - name: openpilot compile3 0.11.0 driving_vision
       run: BENCHMARK_LOG=openpilot_0_11_0_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_vision.onnx
+    - name: openpilot compile3 0.11.0 driving_vision (from pickle)
+      run: BENCHMARK_LOG=openpilot_0_11_0_vision_run_pickle RUN_PICKLE=1 PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM taskset -c 4-7 python3 examples/openpilot/compile3.py
     - name: IR3 openpilot compile3 0.11.0 driving_vision
       run: BENCHMARK_LOG=ir3_openpilot_0_11_0_vision PYTHONPATH="." ASSERT_MIN_STEP_TIME=17 DEV=QCOM:IR3 FLOAT16=1 IMAGE=1 NOLOCALS=1 taskset -c 4-7 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/v0.11.0/selfdrive/modeld/models/driving_vision.onnx
     - name: openpilot compile3 0.11.0 driving_policy
@@ -646,7 +649,7 @@ jobs:
         # benchmark on DSP with NOOPT=1, the devectorizer has issues
         PYTHONPATH=. CC=clang-19 DEV=DSP NOOPT=1 CNT=2 DEBUG=2 python3 examples/test_onnx_imagenet.py /tmp/model.quant.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: ./.github/actions/process-replay
 
   testcommausbgpubenchmark:
     name: UsbGPU Benchmark (comma)
@@ -668,6 +671,8 @@ jobs:
       run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision PYTHONPATH="." GMMU=0 DEV=USB+AMD:LLVM ASSERT_MIN_STEP_TIME=50 python3 examples/openpilot/compile3.py https://github.com/commaai/openpilot/raw/720392c9a5b986981fdbed1bb8c47a6c5573a50e/selfdrive/modeld/models/driving_vision.onnx
     - name: openpilot load_pickle 0.10.1 driving_vision
       run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision_load_pickle PYTHONPATH="." GMMU=0 DEV=USB+AMD ASSERT_MIN_LOAD_TIME=15 python3 examples/openpilot/load_pickle.py
+    - name: openpilot run_pickle 0.10.1 driving_vision
+      run: BENCHMARK_LOG=usbgpu_openpilot_0_10_1_vision_run_pickle RUN_PICKLE=1 PYTHONPATH="." GMMU=0 DEV=USB+AMD ASSERT_MIN_STEP_TIME=50 python3 examples/openpilot/compile3.py
 
   testreddriverbenchmark:
     name: AM Benchmark
@@ -741,7 +746,7 @@ jobs:
         DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=PCI+AMD AMD_AQL=1 python3 test/test_tiny.py
         pkill -f 'extra/remote/serve.py' || true
     - 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: ./.github/actions/process-replay
 
   testgreendriverbenchmark:
     name: NV Benchmark
@@ -804,4 +809,4 @@ jobs:
         DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6483 DEV=NV python3 test/test_tiny.py
         pkill -f 'extra/remote/serve.py' || true
     - 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: ./.github/actions/process-replay

.github/workflows/test.yml

@@ -1,7 +1,7 @@
 name: Unit Tests
 env:
   # increment this when downloads substantially change to avoid the internet
-  CACHE_VERSION: '18'
+  CACHE_VERSION: '19'
   CAPTURE_PROCESS_REPLAY: 1
   GH_TOKEN: ${{ secrets.GITHUB_TOKEN }}
   PYTHONPATH: ${{ github.workspace }}
@@ -333,7 +333,7 @@ jobs:
         deps: testing_unit
         python-version: '3.14'
     - name: Test SPEC=2
-      run: SPEC=2 pytest --maxfail=10 -n auto --durations=30 test/unit test/backend test/opt --ignore test/backend/test_custom_kernel.py --ignore test/unit/test_hashing.py --timeout 60 -k "not test_setitem_big" --splits 2 --group ${{ matrix.group }}
+      run: SPEC=2 pytest --maxfail=10 -n auto --durations=30 test/unit test/backend test/opt --ignore test/backend/test_custom_kernel.py --ignore test/unit/test_hashing.py --timeout 60 -k "not test_setitem_big" -k "not test_conv2d_ceildiv_edge_case" --splits 2 --group ${{ matrix.group }}
 
   fuzzing:
     name: Fuzzing
@@ -417,7 +417,7 @@ jobs:
           llvm: 'true'
       - name: Test openpilot model kernel count and gate usage
         run: |
-          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1486 ALLOWED_GATED_READ_IMAGE=17 FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
+          ALLOWED_KERNEL_COUNT=123 ALLOWED_READ_IMAGE=1486 ALLOWED_GATED_READ_IMAGE=18 FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
       - name: Test openpilot CL compile fp16
         run: FLOAT16=1 DEV=CL IMAGE=1 python examples/openpilot/compile3.py https://gitlab.com/commaai/openpilot-lfs.git/gitlab-lfs/objects/cf6376aa9a090f0da26c280ef69eabf9bbdd51d1faac9ed392919c3db69be916
       - name: Test openpilot CL compile fp32 (test correctness)
@@ -505,14 +505,14 @@ jobs:
         with:
           key: apps_llm
       - name: Test 1B LLM (llama)
-        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model llama3.2:1b | tee /dev/stderr | grep -i rooster
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model llama3.2:1b | tee /dev/stderr | grep -i rooster
       - name: Test 1B LLM (llama q4)
-        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model llama3.2:1b-q4 | tee /dev/stderr | grep -i rooster
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model llama3.2:1b-q4 | tee /dev/stderr | grep -i rooster
       - name: Test 1B LLM (qwen3.5)
-        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model qwen3.5:0.8b | tee /dev/stderr | grep -i rooster
+        run: echo "What's a male chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model qwen3.5:0.8b | tee /dev/stderr | grep -i rooster
       - name: Test 1B LLM (qwen)
         # NOTE: qwen is dumb and only knows about female chickens
-        run: echo "What's a female chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.apps.llm --model qwen3:0.6b | tee /dev/stderr | grep -i hen
+        run: echo "What's a female chicken called? Answer with only one word." | MAX_BUFFER_SIZE=0 python3 -m tinygrad.llm --model qwen3:0.6b | tee /dev/stderr | grep -i hen
 
 # ****** Models Tests ******
 
@@ -643,8 +643,7 @@ jobs:
     runs-on: ubuntu-24.04
     timeout-minutes: 20
     env:
-      DEV: AMD
-      MOCKGPU: 1
+      DEV: MOCKKFD+AMD
     steps:
       - name: Checkout Code
         uses: actions/checkout@v6
@@ -670,29 +669,28 @@ jobs:
       - name: Run AMD renderer tests
         run: python -m pytest -n=auto test/amd/ --durations 20
       - name: Run AMD renderer tests (AMD:LLVM)
-        run: DEV=AMD:LLVM python -m pytest -n=auto test/amd/ --durations 20
+        run: DEV=MOCKKFD+AMD:LLVM python -m pytest -n=auto test/amd/ --durations 20
       - name: Run SQTT profiling tests
         run: PROFILE=1 SQTT=1 python3 -m pytest -n=auto test/amd/test_sqtt_profiler.py
       - name: Run AMD emulated tests on NULL backend
         env:
           AMD: 0
         run: |
-          PYTHONPATH=. DEV=NULL::gfx1100 python extra/mmapeak/mmapeak.py
-          PYTHONPATH=. DEV=NULL::gfx1201 python3 -m pytest -n=auto test/testextra/test_tk.py test/backend/test_asm_gemm.py
-      - name: Run matmul on MOCKGPU
+          PYTHONPATH=. DEV=NULL:HIP:gfx1100 python extra/mmapeak/mmapeak.py
+          PYTHONPATH=. DEV=NULL:HIP:gfx950 python3 -m pytest -n=auto test/testextra/test_tk.py test/backend/test_asm_gemm.py
+      - name: Run matmul on MOCKKFD
         run: |
-          PYTHONPATH="." DEV=AMD MOCKGPU=1 N=256 python3 extra/gemm/amd_asm_matmul.py
-          PYTHONPATH="." DEV=AMD MOCKGPU=1 N=256 python3 extra/gemm/amd_copy_matmul.py
+          PYTHONPATH="." DEV=MOCKKFD+AMD N=256 python3 extra/gemm/amd_asm_matmul.py
+          PYTHONPATH="." DEV=MOCKKFD+AMD N=256 python3 extra/gemm/amd_copy_matmul.py
       - name: Run LLVM test
-        run: DEV=AMD:LLVM python test/device/test_amd_llvm.py
+        run: DEV=MOCKKFD+AMD:LLVM python test/device/test_amd_llvm.py
 
   testmockam:
     name: Linux (am)
     runs-on: ubuntu-24.04
     timeout-minutes: 15
     env:
-      DEV: PCI+AMD
-      MOCKGPU: 1
+      DEV: MOCKPCI+AMD
     steps:
       - name: Checkout Code
         uses: actions/checkout@v6
@@ -704,13 +702,13 @@ jobs:
           amd: 'true'
       - name: Run test_tiny on MOCKAM
         run: python test/test_tiny.py
-      - name: Run test_tiny on MOCKAM USB
-        run: GMMU=0 DEV=USB+AMD python test/test_tiny.py
-      - name: Run test_hcq on MOCKAM
+      - name: Run test_tiny on MOCKUSB
+        run: GMMU=0 DEV=MOCKUSB+AMD python test/test_tiny.py
+      - name: Run test_hcq on MOCKPCI
         run: python -m pytest test/device/test_hcq.py
-      - name: Run disk copy tests on MOCKAM
+      - name: Run disk copy tests on MOCKPCI
         run: python -m pytest test/unit/test_disk_tensor.py -k test_copy_from_disk
-      - name: Run test_tiny on MOCKAM Remote
+      - name: Run test_tiny on MOCKPCI Remote
         run: |
           python extra/remote/serve.py 6667 &
           sleep 2
@@ -728,8 +726,7 @@ jobs:
     runs-on: ubuntu-22.04
     timeout-minutes: 15
     env:
-      DEV: AMD:${{ matrix.backend == 'amdllvm' && 'LLVM' || '' }}:${{ matrix.arch }}
-      MOCKGPU: 1
+      DEV: MOCKKFD+AMD:${{ matrix.backend == 'amdllvm' && 'LLVM' || '' }}:${{ matrix.arch }}
       SKIP_SLOW_TEST: 1
     steps:
       - name: Checkout Code
@@ -764,7 +761,6 @@ jobs:
     runs-on: ubuntu-22.04
     timeout-minutes: 20
     env:
-      MOCKGPU: 1
       FORWARD_ONLY: 1
     steps:
       - name: Checkout Code
@@ -777,7 +773,7 @@ jobs:
           cuda: 'true'
           ocelot: 'true'
       - name: Set env
-        run: printf "${{ matrix.backend == 'ptx' && 'DEV=CUDA:PTX' || matrix.backend == 'nv' && 'DEV=NV\nSKIP_SLOW_TEST=1' }}" >> $GITHUB_ENV
+        run: printf "${{ matrix.backend == 'ptx' && 'DEV=MOCK+CUDA:PTX' || matrix.backend == 'nv' && 'DEV=MOCK+NV\nSKIP_SLOW_TEST=1' }}" >> $GITHUB_ENV
       - name: Check Device.DEFAULT and print some source
         run: |
           python3 -c "from tinygrad import Device; assert Device.DEFAULT in ['CUDA','NV'], Device.DEFAULT"
@@ -862,22 +858,19 @@ jobs:
       run: DEV=METAL TRANSCENDENTAL=2 python -m pytest -n=auto test/backend/test_ops.py::TestOps::test_sin test/backend/test_ops.py::TestOps::test_cos test/backend/test_ops.py::TestOps::test_tan test/backend/test_ops.py::TestOps::test_exp test/backend/test_ops.py::TestOps::test_log --durations=20
     - name: Run pytest (amd)
       env:
-        MOCKGPU: 1
-        DEV: AMD
+        DEV: MOCKKFD+AMD
         FORWARD_ONLY: 1
       run: |
         python3 -m pytest -n=auto test/device/test_hcq.py test/test_tiny.py --durations=20
     - name: Run pytest (amd with llvm backend)
       env:
-        MOCKGPU: 1
-        DEV: "AMD:LLVM"
+        DEV: "MOCKKFD+AMD:LLVM"
         FORWARD_ONLY: 1
       run: |
         python -m pytest -n=auto test/device/test_hcq.py test/test_tiny.py test/device/test_amd_llvm.py --durations=20
     - name: Run pytest (ptx)
       env:
-        MOCKGPU: 1
-        DEV: "NV:PTX"
+        DEV: "MOCK+NV:PTX"
         FORWARD_ONLY: 1
         # TODO: failing due to library loading error
         CAPTURE_PROCESS_REPLAY: 0
@@ -952,8 +945,8 @@ jobs:
       - name: Run process replay tests
         uses: ./.github/actions/process-replay
       - name: Run macOS-specific unit test
-        if: matrix.backend == 'cpu'
-        run: python3 -m pytest test/unit/test_disk_tensor.py::TestDiskTensor::test_copy_to_cpu_not_truncated
+        if: matrix.backend == 'llvm'
+        run: python3 -m pytest test/unit/test_disk_tensor.py::TestDiskTensor::test_copy_to_cpu_not_truncated test/unit/test_cpu.py
 
 # ****** Windows Tests ******
 

.gitignore

@@ -68,3 +68,4 @@ mutants
 .mutmut-cache
 dagre/
 graphlib/
+uv.lock

docs/abstractions3.py

@@ -1,6 +1,4 @@
 # abstractions2 goes from back to front, here we will go from front to back
-from typing import List
-from tinygrad.helpers import tqdm
 
 # *****
 # 0. Load mnist on the device
@@ -33,21 +31,21 @@ model(X).sparse_categorical_crossentropy(Y).backward()
 optim.schedule_step()   # this will step the optimizer without running realize
 
 # *****
-# 3. Create a schedule.
+# 3. Create a schedule (linear uop).
 
 # The weight Tensors have been assigned to, but not yet realized. Everything is still lazy at this point
 # l1.uop and l2.uop define a computation graph
 
-from tinygrad.schedule import ExecItem
-schedule: List[ExecItem] = Tensor.schedule(l1, l2)
+from tinygrad.engine.realize import run_linear
+linear = Tensor.schedule_linear(l1, l2)
 
-print(f"The schedule contains {len(schedule)} items.")
-for si in schedule: print(str(si)[:80])
+print(f"The schedule contains {len(linear.src)} items.")
+for call in linear.src: print(str(call)[:80])
 
 # *****
-# 4. Lower and run the schedule.
+# 4. Lower and run the schedule (linear uop).
 
-for si in tqdm(schedule): si.run()
+run_linear(linear)
 
 # *****
 # 5. Print the weight change

docs/abstractions4.py

@@ -1,9 +1,9 @@
 # tinygrad allows you to write kernels at many different abstractions levels.
 # This is for RDNA3, but if you don't have one you can run with the emulator
-# PYTHONPATH="." MOCKGPU=1 DEV=AMD
+# PYTHONPATH="." DEV=MOCKPCI+AMD
 
 from tinygrad import Tensor, Context, GlobalCounters, UOp, Device
-from tinygrad.helpers import DEBUG, getenv
+from tinygrad.helpers import DEV, DEBUG, getenv
 from tinygrad.uop.ops import AxisType, KernelInfo, Ops
 from tinygrad.dtype import AddrSpace, dtypes
 from tinygrad.runtime.autogen.amd.rdna3.ins import *
@@ -16,7 +16,7 @@ def eval_harness(name, tensor, fxn, check=None):
   print(f"computed in {GlobalCounters.time_sum_s*1000:.2f} ms, {(a.nbytes()/1e9)/GlobalCounters.time_sum_s:.2f} GB/s")
   return out
 
-SZ = 256*1024 if getenv("MOCKGPU") else 1024*1024*1024
+SZ = 256*1024 if DEV.interface.startswith("MOCK") else 1024*1024*1024
 
 def example_2_hip(a:Tensor, correct):
   GLOBALS = 1024
@@ -105,7 +105,7 @@ def example_3_custom_uop(a:Tensor, correct):
 def example_5_custom_assembly(a:Tensor, correct):
   # Kernel class copied from amd_asm_matmul
   class Kernel:
-    def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
+    def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
     def label(self, name): self.labels[name] = self.pos
     def emit(self, inst, target=None):
       self.instructions.append(inst)

docs/developer/developer.md

@@ -17,15 +17,13 @@ The `UOp` graph specifies the compute in terms of low level tinygrad ops. Not al
 
 ## Scheduling
 
-The [scheduler](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/schedule/__init__.py) converts the graph of UOps into a list of `ExecItem`. One `ExecItem` is one kernel on the GPU, and the scheduler is responsible for breaking the large compute graph into subgraphs that can fit in a kernel. `ast` specifies what compute to run, and `bufs` specifies what buffers to run it on.
-
-::: tinygrad.schedule.ExecItem
+The [scheduler](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/schedule/__init__.py) converts the graph of UOps into a `LINEAR` UOp whose `src` is a list of `CALL` UOps. One `CALL` is one kernel on the GPU, and the scheduler is responsible for breaking the large compute graph into subgraphs that can fit in a kernel. The `CALL`'s `src[0]` (a `SINK` ast) specifies what compute to run, and the remaining `src` are the buffers to run it on.
 
 ## Lowering
 
-The code in [realize](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/engine/realize.py) lowers `ExecItem` by populating its `prg` field with
+The code in [realize](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/engine/realize.py) lowers each `CALL` by compiling its ast into a `PROGRAM` and running it.
 
-::: tinygrad.engine.realize.run_schedule
+::: tinygrad.engine.realize.run_linear
 
 There's a ton of complexity hidden behind this, see the `codegen/` directory.
 
@@ -35,13 +33,7 @@ Then we render the UOps into code with a `Renderer`, then we compile the code to
 
 ## Execution
 
-Creating `ExecItem`, which has a run method
-
-::: tinygrad.engine.realize.ExecItem
-    options:
-        members: true
-
-Lists of `ExecItem` can be condensed into a single ExecItem with the Graph API (rename to Queue?)
+`run_linear` walks the `LINEAR` UOp, dispatching each `CALL` to a runner (kernel, copy, view, encdec, or graph).
 
 ## Runtime
 

docs/developer/layout.md

@@ -28,7 +28,7 @@ Transforms the ast into an optimized ast. This is where BEAM search and heuristi
 
 Transform the optimized ast into a linearized and rendered program.
 
-::: tinygrad.codegen.get_program
+::: tinygrad.codegen.to_program
     options:
         members: false
         show_labels: false
@@ -53,7 +53,7 @@ Transform the linearized list of UOps into a program, represented as a string.
 
 Abstracted high level interface to the runtimes.
 
-::: tinygrad.engine.realize.get_program
+::: tinygrad.engine.realize.to_program
     options:
         members: false
         show_labels: false

docs/env_vars.md

@@ -57,6 +57,8 @@ AMD:LLVM      | use the AMD device with the LLVM renderer
 NV:CUDA:sm_70 | use the NV device with the CUDA renderer targetting sm_70
 AMD::gfx950   | use the AMD device targetting gfx950
 USB+AMD       | use the AMD device over the USB interface
+CPU:LLVM      | use the CPU device with the LLVM renderer
+CPU:LLVM:x86_64,znver2,avx2,-avx512f | use the CPU device with the LLVM renderer, with [additional arch flags](runtime.md#cpu-arch)
 
 ### Debug breakdown
 

docs/nn.md

@@ -37,4 +37,4 @@
     options:
         show_signature: false
         separate_signature: false
-::: tinygrad.nn.state.gguf_load
+::: tinygrad.llm.gguf.gguf_load

docs/runtime.md

@@ -5,12 +5,12 @@ tinygrad supports various runtimes, enabling your code to scale across a wide ra
 | Runtime | Description | Compiler Options | Requirements |
 |---------|-------------|------------------|--------------|
 | [NV](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_nv.py) | Provides acceleration for NVIDIA GPUs | nvrtc (default)<br>PTX (`DEV=NV:PTX`) | Ampere/Ada/Blackwell series GPUs.<br>You can select an interface via [the `DEV` variable](env_vars.md#dev-variable). See [NV interfaces](#nv-interfaces) for details. |
-| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | LLVM (`DEV=AMD:LLVM`)<br>HIP/COMGR (`DEV=AMD:HIP`) | RDNA2 or newer GPUs.<br>You can select an interface via [the `DEV` variable](env_vars.md#dev-variable). See [AMD interfaces](#amd-interfaces) for details. |
+| [AMD](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_amd.py) | Provides acceleration for AMD GPUs | LLVM (`DEV=AMD:LLVM`)<br>HIP/COMGR (`DEV=AMD:HIP`) | CDNA3, CDNA4, RDNA3 or RDNA4 GPUs.<br>You can select an interface via [the `DEV` variable](env_vars.md#dev-variable). See [AMD interfaces](#amd-interfaces) for details. |
 | [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 | nvrtc (default)<br> PTX (`DEV=CUDA:PTX`) | NVIDIA GPU with CUDA support |
 | [CL](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cl.py) | Accelerates computations using OpenCL on GPUs | - | OpenCL 2.0 compatible device |
-| [CPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cpu.py) | Runs on CPU using the clang or llvm compiler | Clang JIT (default)<br>LLVM IR (`DEV=CPU:LLVM`) | `clang` compiler in system `PATH` |
+| [CPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_cpu.py) | Runs on CPU using the clang or llvm compiler | Clang JIT (default)<br>LLVM IR (`DEV=CPU:LLVM`) | `clang` compiler in system `PATH`<br>You can specify additional arch parameters via [the `DEV` variable](env_vars.md#dev-variable). See [CPU arch](#cpu-arch) for details. |
 | [WEBGPU](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/ops_webgpu.py) | Runs on GPU using the Dawn WebGPU engine (used in Google Chrome) | - | Dawn library installed and discoverable. Binaries: [pydawn v0.3.0](https://github.com/wpmed92/pydawn/releases/tag/v0.3.0) |
 
 
@@ -79,3 +79,13 @@ NV backend supports several interfaces for communicating with devices:
 
 * `NVK`: uses the nvidia driver
 * `PCI`: uses the [NV driver](https://github.com/tinygrad/tinygrad/tree/master/tinygrad/runtime/support/nv/nvdev.py)
+
+## CPU Arch
+The CPU renderers may be additionally configured using the arch component of [the `DEV` environment variable](env_vars.md#dev-variable).
+CPU arch should be specified as a comma-separated list of parameters, and must contain at least two values: the architecture family (ie. x86_64, arm64, or riscv64) and the cpu type (as accepted by `clang`'s `-march`).
+If native is specified as the cpu type, tinygrad (or delegate compiler) will query the host cpu type. Additional comma-separated values may be specified as follows:
+
+* `AMX`: emit Apple silicon AMX instructions
+
+All other additional values are interpreted as cpu feature flags. When a value is preceded by a `-` character, the corresponding feature flag will be disabled, otherwise the flag will be enabled.
+Note that enabled feature flags should not be preceded by a `+`.

docs/tensor/elementwise.md

@@ -66,8 +66,8 @@ Elementwise ops operate on a per element basis. They don't change the shape of t
 ::: tinygrad.Tensor.sub
 ::: tinygrad.Tensor.mul
 ::: tinygrad.Tensor.div
-::: tinygrad.Tensor.idiv
 ::: tinygrad.Tensor.mod
+::: tinygrad.Tensor.fmod
 ::: tinygrad.Tensor.bitwise_xor
 ::: tinygrad.Tensor.bitwise_and
 ::: tinygrad.Tensor.bitwise_or

docs/tensor/properties.md

@@ -19,8 +19,8 @@
 
 ## tinygrad ops
 
-::: tinygrad.Tensor.schedule_with_vars
-::: tinygrad.Tensor.schedule
+::: tinygrad.Tensor.linear_with_vars
+::: tinygrad.Tensor.schedule_linear
 ::: tinygrad.Tensor.realize
 ::: tinygrad.Tensor.replace
 ::: tinygrad.Tensor.assign

docs/tinygpu.md

@@ -4,7 +4,7 @@ TinyGPU app lets you use AMD and NVIDIA GPUs on macOS over USB4/Thunderbolt with
 
 ## Requirements
 
-- macOS (12.1+)
+- macOS (13.0+)
 - USB4/Thunderbolt port
 - A supported GPU (AMD RDNA3+ or NVIDIA Ampere+)
 
@@ -55,7 +55,7 @@ export PATH="$HOME/.local/bin:$PATH"
 ### 5. Use it!
 
 ```bash
-DEV={AMD|NV} python3 tinygrad/apps/llm.py
+DEV={AMD|NV} python3 -m tinygrad.llm
 ```
 
 **Note:** Use `JITBEAM=2` to search for faster kernels (one-time search cost, results cached).

examples/anthropic_challenge.py

@@ -113,7 +113,7 @@ class VLIWRenderer(Renderer):
         case Ops.GEP:
           # a GEP is just an alias to a special register in the vector
           r[u] = r[u.src[0]] + u.arg[0]
-        case Ops.VECTORIZE:
+        case Ops.STACK:
           if all(s == u.src[0] for s in u.src):
             # if all sources are the same, we can broadcast
             inst.append({"valu": [("vbroadcast", r[u], r[u.src[0]])]})
@@ -173,16 +173,16 @@ if __name__ == "__main__":
 
   # *** render to device ***
 
-  from tinygrad.codegen import get_program
+  from tinygrad.codegen import to_program
   with Context(PCONTIG=2, DEVECTORIZE=2, SPEC=0):
     out = tree_traversal(forest_t, val_t, height, rounds)
-    sink = out.schedule()[-1].ast
-    prg = get_program(sink, VLIWRenderer())
+    sink = out.schedule_linear().src[-1].src[0]
+    prg = to_program(sink, VLIWRenderer())
 
   # *** run on Machine and compare ***
 
   # NOTE: the scratch size needs to be reduced to 1536 when you have a register allocator
-  src = eval(prg.src)
+  src = eval(prg.src[3].arg)
   max_regs = max(t[1] for instr in src for v in instr.values() for t in v if len(t) > 1) + 8
   print(f"{max_regs:5d} regs used" + ("" if max_regs <= 1536 else "       <-- WARNING: TOO MANY REGISTERS, MUST BE <= 1536"))
   machine = problem.Machine(mem, src, problem.DebugInfo(scratch_map={}), n_cores=1, trace=False, scratch_size=max_regs)

examples/compile_tensorflow.py

@@ -35,12 +35,11 @@ def compile_onnx_model(onnx_model):
   tinyonnx = TinyOnnx(onnx_model)
   the_input = Tensor.randn(1,32)
 
-  run, special_names = jit_model(tinyonnx, the_input)
+  linear, output_bufs = jit_model(tinyonnx, the_input)
+  the_output = [tinyonnx.forward(the_input)]
 
-  functions, statements, bufs, bufs_to_save = compile_net(run, special_names)
+  functions, statements, bufs, bufs_to_save = compile_net(linear, output_bufs)
   prg = export_model_clang(functions, statements, bufs, {}, ["input0"], ["output0"])
-
-  the_output = run(the_input)
   cprog = ["#include <string.h>", "#include <stdio.h>", "#include <stdlib.h>"]
   cprog.append(prg)
 

examples/gpt2.py

@@ -5,8 +5,9 @@ with contextlib.suppress(ImportError): import tiktoken
 from tinygrad import Tensor, TinyJit, Device, GlobalCounters, Variable, dtypes
 from tinygrad.uop.ops import UOp
 from tinygrad.helpers import Timing, DEBUG, JIT, getenv, fetch, colored, trange
+from tinygrad.llm.gguf import gguf_load
 from tinygrad.nn import Embedding, Linear, LayerNorm
-from tinygrad.nn.state import gguf_load, torch_load, load_state_dict, get_state_dict
+from tinygrad.nn.state import torch_load, load_state_dict, get_state_dict
 from extra.bench_log import BenchEvent, WallTimeEvent
 
 MAX_CONTEXT = getenv("MAX_CONTEXT", 128)

examples/llama3.py

@@ -2,7 +2,8 @@ from pathlib import Path
 from typing import List
 import json, argparse, random, time, os
 from extra.models.llama import Transformer, convert_from_huggingface, convert_from_gguf, fix_bf16
-from tinygrad.nn.state import safe_load, torch_load, load_state_dict, get_parameters, gguf_load
+from tinygrad.llm.gguf import gguf_load
+from tinygrad.nn.state import safe_load, torch_load, load_state_dict, get_parameters
 from tinygrad import Tensor, dtypes, nn, Context, Device, GlobalCounters
 from tinygrad.helpers import Profiling, Timing, DEBUG, colored, fetch, tqdm
 from extra.bench_log import BenchEvent, WallTimeEvent
@@ -122,7 +123,7 @@ def NF4Linear(block_size):
     def __call__(self, x: Tensor) -> Tensor:
       high_bits = self.weight
       low_bits = (self.weight * 2 ** 4).contiguous()
-      unpacked = Tensor.stack(high_bits, low_bits, dim=-1).idiv(2 ** 4)
+      unpacked = Tensor.stack(high_bits, low_bits, dim=-1).div(2 ** 4, rounding_mode="trunc")
       unscaled = CODE[unpacked].to(x.device).reshape(-1, block_size) * self.scale
       return x.linear(unscaled.reshape(self.out_features, self.in_features).T)
 

examples/mlperf/model_train.py

@@ -1282,7 +1282,7 @@ def train_bert():
         previous_step = i
 
 def train_llama3():
-  from examples.mlperf.models.flat_llama import FlatTransformer, apply_grad, FP8
+  from examples.mlperf.models.flat_llama import FlatTransformer, apply_grad, FP8_DTYPE
   from examples.llama3 import MODEL_PARAMS
   from examples.mlperf.lr_schedulers import CosineAnnealingLRWithWarmup
   from examples.mlperf.optim import GradAccClipAdamW
@@ -1357,6 +1357,7 @@ def train_llama3():
       MLLOGGER.event(key=mllog_constants.OPT_LR_WARMUP_STEPS, value=WARMUP_STEPS)
       MLLOGGER.event(key=mllog_constants.NUM_WARMUP_STEPS, value=WARMUP_STEPS)
       MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_STEPS, value=MAX_STEPS - WARMUP_STEPS)
+      MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_SCHEDULE, value="cosine with linear warmup")
       MLLOGGER.event(key=mllog_constants.OPT_GRADIENT_CLIP_NORM, value=1.0)
   else:
     MLLOGGER = None
@@ -1395,7 +1396,7 @@ def train_llama3():
 
   params = get_parameters(model)
 
-  if getenv("FAKEDATA"):
+  if getenv("EMPTYWEIGHT"):
     for v in get_parameters(model):
       v = v.assign(Tensor.empty(v.shape, dtype=v.dtype))
 
@@ -1416,9 +1417,12 @@ def train_llama3():
   optim = GradAccClipAdamW(params, lr=0.0, b1=opt_adamw_beta_1, b2=opt_adamw_beta_2,
                            eps=opt_adamw_epsilon, weight_decay=opt_adamw_weight_decay, grad_acc=grad_acc, device=optim_device)
 
-  # init grads
   for p in optim.params:
-    p.grad = Tensor.zeros(p.shape, dtype=p.dtype, device=p.device).contiguous()
+    grad_dtype = dtypes.bfloat16 if p.dtype == FP8_DTYPE else p.dtype
+    if isinstance(p.device, tuple) and p.uop.axis is not None:
+      p.grad = Tensor.zeros(p.shape, dtype=grad_dtype, device=p.device[0]).shard_(p.device, axis=p.uop.axis).contiguous()
+    else:
+      p.grad = Tensor.zeros(p.shape, dtype=grad_dtype, device=p.device).contiguous()
   grads = [p.grad for p in optim.params]
 
   scheduler = CosineAnnealingLRWithWarmup(optim, opt_base_learning_rate, opt_end_learning_rate, opt_learning_rate_warmup_steps, opt_learning_rate_decay_steps)
@@ -1432,7 +1436,22 @@ def train_llama3():
     print(f"loading optim checkpoint from {fn}")
     load_state_dict(scheduler, safe_load(fn), realize=False)
 
-  fp8_amax = [t for ts in model._fp8_amax.values() for t in ts] if FP8 else []
+  fp8_amax = [t for ts in model._fp8_amax.values() for t in ts]
+  fp8_grad_amax = [t for ts in model._fp8_grad_amax.values() for t in ts] if hasattr(model, "_fp8_grad_amax") else []
+  fp8_inv_scales = list(model._fp8_inv_scale.values())
+
+  from tinygrad.nn.state import get_state_dict
+  model_state = get_state_dict(model)
+  for wname in ["wqkv", "wo", "w13", "w2"]:
+    w = model_state[wname]
+    w._inv_scale = model._fp8_inv_scale[wname]
+    if optim.master_params:
+      idx = next(j for j, p in enumerate(optim.params) if p is w)
+      optim.master_params[idx].assign((optim.master_params[idx] * w._inv_scale.reshape(-1, *([1]*(w.ndim-1)))).contiguous())
+
+  # realize everything here
+  if optim.master_params: Tensor.realize(*optim.master_params)
+  Tensor.realize(*optim.params, *fp8_inv_scales, *fp8_amax, *fp8_grad_amax)
 
   @TinyJit
   def minibatch(tokens:Tensor):
@@ -1440,13 +1459,17 @@ def train_llama3():
     if is_mp: tokens = tokens.shard(device)
     if not is_sharding: tokens = tokens.to(None)
     logits:Tensor = model(tokens[:, :-1])
-    loss = vocab_mask.where(-1e9, logits).sparse_categorical_crossentropy(tokens[:, 1:])
+    if getenv("FAST_CE", 0):
+      from extra.llama_kernels.fused_ce import fused_ce_loss
+      loss = fused_ce_loss(logits.cast(dtypes.bfloat16), tokens[:, 1:], label_smoothing=0.0)
+    else:
+      loss = vocab_mask.where(-1e9, logits).sparse_categorical_crossentropy(tokens[:, 1:])
 
     for g, new_g in zip(grads, loss.gradient(*optim.params)):
       apply_grad(g, new_g.uop)
 
     loss_cpu = loss.flatten().float().to("CPU")
-    return loss_cpu.realize(*grads, *fp8_amax)
+    return loss_cpu.realize(*grads, *fp8_amax, *fp8_grad_amax)
 
   @TinyJit
   def optim_step():
@@ -1457,7 +1480,7 @@ def train_llama3():
 
     lr_cpu = optim.lr.float().to("CPU")
     grad_norm_cpu = grad_norm.float().to("CPU")
-    Tensor.realize(lr_cpu, grad_norm_cpu, *grads)
+    Tensor.realize(lr_cpu, grad_norm_cpu, *grads, *fp8_inv_scales)
 
     return lr_cpu, grad_norm_cpu
 
@@ -1544,7 +1567,7 @@ def train_llama3():
 
       mem_gb = GlobalCounters.mem_used / 1e9
       gflops = GlobalCounters.global_ops / 1e9 / dev_time
-      mfu = ((6 * num_params * SEQLEN * GBS) / (dev_time * device_count * (4.6e15 if FP8 else 2.3e15))) * 100
+      mfu = ((6 * num_params * SEQLEN * GBS) / (dev_time * device_count * 4.6e15)) * 100
       tqdm.write(
           f"{i:5} {step_time:.3f} s step, {gbs_time:.3f} s gbs, {optim_time:.3f} s optim, {data_time:.3f} s data, {loss:.4f} loss, " \
           f"{lr:.12f} LR, {grad_norm:.6f} grad_norm, {mem_gb:.2f} GB used, {gflops:9.2f} GFLOPS, {mfu:5.2f}% MFU")
@@ -1621,7 +1644,6 @@ def train_llama3():
         tqdm.write(f"target achieved after {sequences_seen} sequences")
         if MLLOGGER and RUNMLPERF:
           MLLOGGER.end(key=mllog_constants.EPOCH_STOP, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
-          MLLOGGER.event(key=mllog_constants.TRAIN_SAMPLES, value=sequences_seen)
           MLLOGGER.end(key=mllog_constants.RUN_STOP, metadata={mllog_constants.STATUS: mllog_constants.SUCCESS})
         if getenv("CKPT"):
           if not os.path.exists(ckpt_dir := "./ckpts"): os.mkdir(ckpt_dir)

examples/mlperf/models/flat_llama.py

@@ -1,4 +1,4 @@
-import math, os, functools
+import math, os
 if __name__ == "__main__":
   os.environ["DEFAULT_FLOAT"] = "bfloat16"
   os.environ["OPTIM_DTYPE"] = "bfloat16"
@@ -16,65 +16,78 @@ from tinygrad import Tensor, nn, function, getenv, dtypes, TinyJit
 from tinygrad.helpers import Timing, colored, GlobalCounters, profile_marker
 from tinygrad.uop.ops import Ops, UOp
 from extra.models.llama import apply_rotary_emb, precompute_freqs_cis
+from extra.llama_kernels.rmsnorm import rmsnorm
+from extra.llama_kernels import FP8_MAX, local_abs_max
 
-FP8 = getenv("FP8", 0)
+ASM_GEMM = getenv("ASM_GEMM", 0)
+FUSED_INPUT_QUANTIZE = getenv("FUSED_INPUT_QUANTIZE", 0)
+FUSED_ADD_NORM_MUL_QUANTIZE = getenv("FUSED_ADD_NORM_MUL_QUANTIZE", 0)
+FUSED_SILU_W13 = getenv("FUSED_SILU_W13", 0)
 
 FP8_DTYPE = dtypes.fp8e4m3
 FP8_GRAD_DTYPE = dtypes.fp8e5m2
-FP8_MAX = 448.0
-
-# per-device abs max without allreduce (matches TE delayed scaling behavior)
-@functools.cache
-def _local_abs_max_fxn(x_p, device):
-  x = Tensor(x_p, device=device)
-  inner = Tensor(x.uop.src[0]) if x.uop.op is Ops.MULTI else x
-  return (inner.abs().max(),)
-
-def _local_abs_max(x:Tensor) -> Tensor:
-  param = x.as_param(0)
-  fxn = _local_abs_max_fxn(param.uop, x.device)
-  return Tensor(fxn[0].uop.call(x.uop).gettuple(0))
 
 def quantize_fp8(x:Tensor, amax_state:Tensor|None=None):
-  new_amax = (_local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach()
+  new_amax = (local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach().cast(dtypes.float32)
   scale = FP8_MAX / ((amax_state if amax_state is not None else new_amax) + 1e-8)
   x_scaled = x * scale
   x_clamped = x_scaled + (x_scaled.detach().clamp(-FP8_MAX, FP8_MAX) - x_scaled.detach())  # STE
   return x_clamped.cast(FP8_DTYPE), scale.float().reciprocal(), new_amax
 
-def matmul(x:Tensor, w:Tensor, fp8=FP8, amax_x:Tensor|None=None, amax_w:Tensor|None=None) -> tuple[Tensor,...]:
+def matmul(x:Tensor, w:Tensor, fp8:bool=True, amax_x:Tensor|None=None, w_inv_scale:Tensor|None=None,
+           x_fp8:Tensor|None=None, x_scale:Tensor|None=None, x_new_amax:Tensor|None=None,
+           grad_amax_state:Tensor|None=None) -> tuple[Tensor,...]:
   if not fp8:
-    if getenv("ASM_GEMM"):
+    if ASM_GEMM:
       from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
       if can_use_asm_gemm(x, w.T): return (asm_gemm(x, w.T),)
     return (x @ w.T,)
-  x_fp8, x_scale, x_new_amax = quantize_fp8(x, amax_state=amax_x)
-  w_fp8, w_scale, w_new_amax = quantize_fp8(w, amax_state=amax_w)
-  combined_scale = x_scale * w_scale
-  if getenv("ASM_GEMM"):
+  assert w_inv_scale is not None, "fp8 matmul requires w_inv_scale (weights must be stored in fp8 with per-tensor scale)"
+  if x_fp8 is None:
+    if FUSED_INPUT_QUANTIZE and amax_x is not None:
+      from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+      x_fp8, x_scale, x_new_amax, _ = quantize_fp8_delayed(x, amax_x, FP8_DTYPE)
+    else:
+      x_fp8, x_scale, x_new_amax = quantize_fp8(x, amax_state=amax_x)
+  if ASM_GEMM:
     from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
-    if can_use_asm_gemm(x_fp8, w_fp8.T): return asm_gemm(x_fp8, w_fp8.T, combined_scale=combined_scale), x_new_amax, w_new_amax, x_fp8, w_fp8
-  return x_fp8.dot(w_fp8.T, dtype=dtypes.float) * combined_scale, x_new_amax, w_new_amax, x_fp8, w_fp8
+    if can_use_asm_gemm(x_fp8, w.T):
+      return asm_gemm(x_fp8, w.T, x_scale=x_scale, w_scale=w_inv_scale, grad_amax_state=grad_amax_state), x_new_amax, x_fp8, w
+  return x_fp8.dot(w.T, dtype=dtypes.float) * x_scale * w_inv_scale, x_new_amax, x_fp8, w
 
-def _rmsnorm_fwd(x_in:Tensor, eps:float) -> tuple[Tensor, Tensor]:
-  x = x_in.float()
-  rrms = (x.square().mean(-1, keepdim=True) + eps).rsqrt()
-  return (x * rrms).cast(x_in.dtype), rrms
+def norm_quantize_matmul(x:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor, grad_amax_state:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, x_normed, rrms = fused_rmsnorm_mul_quantize_fp8(x, norm, amax_x, eps, FP8_DTYPE)
+    out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax, grad_amax_state=grad_amax_state)
+    return out, x_normed, rrms, ret
+  x_normed, rrms = rmsnorm(x, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale, grad_amax_state=grad_amax_state)
+  return out, x_normed, rrms, ret
 
-@functools.cache
-def _rmsnorm_fwd_fxn(x_in_p, eps, device):
-  return _rmsnorm_fwd(Tensor(x_in_p, device=device), eps)
+def add_norm_quantize_matmul(x:Tensor, residual:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_add_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, h, x_normed, rrms = fused_add_rmsnorm_mul_quantize_fp8(x, residual, norm, amax_x, eps, FP8_DTYPE)
+    out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax)
+    return out, h, x_normed, rrms, ret
+  h = x + residual
+  x_normed, rrms = rmsnorm(h, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale)
+  return out, h, x_normed, rrms, ret
 
-def _rmsnorm_bwd(grad:UOp, call:UOp) -> tuple:
-  x_normed = Tensor(call.gettuple(0)).float()
-  do_float = Tensor(grad).float()
-  d_x = Tensor(call.gettuple(1)) * (do_float - x_normed * (do_float * x_normed).mean(-1, keepdim=True))
-  return (d_x.cast(call.src[1].dtype).uop,)
-
-def rmsnorm(x_in:Tensor, eps:float) -> tuple[Tensor, Tensor]:
-  fxn = _rmsnorm_fwd_fxn(x_in.as_param(0).uop, eps, x_in.device)
-  call = UOp.maketuple(fxn[0].uop, fxn[1].uop).call(x_in.uop, grad_fxn=_rmsnorm_bwd)
-  return Tensor(call.gettuple(0)), Tensor(call.gettuple(1))
+def silu_w13_quantize_matmul(x_w13:Tensor, w2:Tensor, s_2:Tensor,
+                             amax_x2:Tensor,
+                             grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
+  if FUSED_SILU_W13:
+    from extra.llama_kernels.cast_amax import fused_quantize_fp8_w13
+    x2_fp8, x2_inv_scale, new_amax_x2 = fused_quantize_fp8_w13(x_w13, amax_x2, FP8_DTYPE, grad_amax_state=grad_amax_xw13)
+    out, *ret = matmul(None, w2, w_inv_scale=s_2, x_fp8=x2_fp8, x_scale=x2_inv_scale, x_new_amax=new_amax_x2, grad_amax_state=grad_amax_xout)
+    return out, ret
+  hidden = x_w13.shape[-1] // 2
+  x_w1, x_w3 = x_w13[..., :hidden], x_w13[..., hidden:]
+  out, *ret = matmul(x_w1.silu() * x_w3, w2, amax_x=amax_x2, w_inv_scale=s_2, grad_amax_state=grad_amax_xout)
+  return out, ret
 
 class FlatTransformer:
   def __init__(self, dim:int, hidden_dim:int, n_heads:int, n_layers:int, norm_eps:float, vocab_size:int, n_kv_heads:int|None=None,
@@ -85,17 +98,18 @@ class FlatTransformer:
     self.n_kv_heads = n_kv_heads if n_kv_heads is not None else n_heads # n_kv_heads != n_heads implies MQA [arxiv/2307.09288, A.2.1]
     self.head_dim = dim // n_heads
     self.n_rep = self.n_heads // self.n_kv_heads
+    self.hidden_dim = hidden_dim
 
     scaled_std = 0.02 / math.sqrt(2 * n_layers)
 
     # Attention
+    self._init_inv_scales = []  # populated by lin_per_layer
     self.wqkv = self.lin_per_layer(dim, self.n_heads * self.head_dim + self.n_kv_heads * self.head_dim * 2)
     self.wo = self.lin_per_layer(self.n_heads * self.head_dim, dim, std=scaled_std)
 
     # FeedForward
-    self.w1 = self.lin_per_layer(dim, hidden_dim)
+    self.w13 = self.lin_per_layer(dim, hidden_dim * 2)
     self.w2 = self.lin_per_layer(hidden_dim, dim, std=scaled_std)
-    self.w3 = self.lin_per_layer(dim, hidden_dim)
 
     self.norm_eps = norm_eps
     self.attention_norm = Tensor.ones(n_layers, dim).contiguous()
@@ -108,83 +122,89 @@ class FlatTransformer:
     self.output = Tensor.normal(1, vocab_size, dim, mean=0.0, std=0.02, dtype=dtypes.bfloat16)
     self.freqs_cis = precompute_freqs_cis(dim // n_heads, max_context * 2, rope_theta).contiguous().requires_grad_(False)
 
-    if FP8:
-      def _amax(): return Tensor.full((), FP8_MAX).contiguous().requires_grad_(False)
-      names = ["xqkv", "wqkv", "xo", "wo", "x1", "w1", "x2", "w2", "x3", "w3"]
-      # _fp8_amax[name][layer_idx] = scalar amax tensor
-      self._fp8_amax = {name: [_amax() for _ in range(n_layers)] for name in names}
-      self._fp8_amax["xout"] = [_amax()]
-      self._fp8_amax["wout"] = [_amax()]
+    def _amax(): return Tensor.full((), FP8_MAX, dtype=dtypes.float32).contiguous().requires_grad_(False)
+    names = ["xqkv", "xo", "x13", "x2"]
+    self._fp8_amax = {name: [_amax() for _ in range(n_layers)] for name in names}
+    grad_names = ["xqkv", "xo", "xw13", "xout"]
+    self._fp8_grad_amax = {name: [_amax() for _ in range(n_layers)] for name in grad_names}
+    w_names = ["wqkv", "wo", "w13", "w2"]
+    self._fp8_inv_scale = {wname: inv_scales.float().contiguous().requires_grad_(False)
+                           for wname, inv_scales in zip(w_names, self._init_inv_scales)}
+    del self._init_inv_scales
 
   def lin_per_layer(self, in_features:int, out_features:int, std:float=0.02):
-    if getenv("ZEROS"): return Tensor.zeros(self.n_layers, out_features, in_features)
-    return Tensor.normal(self.n_layers, out_features, in_features, mean=0.0, std=std)
+    if getenv("ZEROS"): w = Tensor.zeros(self.n_layers, out_features, in_features)
+    else: w = Tensor.normal(self.n_layers, out_features, in_features, mean=0.0, std=std)
+    amax = w.abs().flatten(1).max(1).detach()
+    scale = FP8_MAX / (amax + 1e-8)
+    self._init_inv_scales.append((amax + 1e-8) / FP8_MAX)
+    return (w * scale.reshape(-1, 1, 1)).clamp(-FP8_MAX, FP8_MAX).cast(FP8_DTYPE)
 
   def attention(self, x:Tensor, freqs_cis:Tensor, attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
-                amax_xqkv=None, amax_wqkv=None, amax_xo=None, amax_wo=None):
+                amax_xqkv:Tensor, amax_xo:Tensor, s_qkv:Tensor, s_o:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor):
     bsz, seqlen, _ = x.shape
     new_amaxs, saves = [], []
 
-    x, rrms = rmsnorm(x, self.norm_eps)
-    saves.extend([x, rrms])
-    x = x * attention_norm
-
-    xqkv, *ret = matmul(x, wqkv, amax_x=amax_xqkv, amax_w=amax_wqkv)
-    new_amaxs.extend(ret[:2])
-    saves.extend(ret[2:] + [xqkv])
+    xqkv, x_normed, rrms, ret = norm_quantize_matmul(x, attention_norm, wqkv, s_qkv, self.norm_eps,
+                                                     amax_x=amax_xqkv, grad_amax_state=grad_amax_xqkv)
+    saves.extend([x_normed, rrms])
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [xqkv])
     xqkv = xqkv.reshape(bsz, seqlen, self.n_kv_heads, self.n_rep + 2, self.head_dim)
     xq = xqkv[:, :, :, :self.n_rep].reshape(bsz, seqlen, self.n_heads, self.head_dim)
     xk = xqkv[:, :, :, self.n_rep].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
     xv = xqkv[:, :, :, self.n_rep+1].reshape(bsz, seqlen, self.n_kv_heads, self.head_dim)
 
     xq, xk = apply_rotary_emb(xq, xk, freqs_cis)
-    if FP8: xq, xk, xv = xq.cast(dtypes.bfloat16), xk.cast(dtypes.bfloat16), xv.cast(dtypes.bfloat16)
-    xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
+    xq, xk, xv = xq.cast(dtypes.bfloat16), xk.cast(dtypes.bfloat16), xv.cast(dtypes.bfloat16)
     if getenv("HK_FLASH_ATTENTION"):
       from extra.thunder.amd.fa import flash_attention
       attn, *save = flash_attention(xq, xk, xv, is_causal=True)
       saves.extend(save)
     else:
-      attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True)
-    attn = attn.transpose(1, 2).reshape(bsz, seqlen, -1)
+      xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
+      attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True).transpose(1, 2)
+    attn = attn.reshape(bsz, seqlen, -1)
 
-    out, *ret = matmul(attn, wo, amax_x=amax_xo, amax_w=amax_wo)
-    new_amaxs.extend(ret[:2])
-    saves.extend(ret[2:] + [out])
+    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o, grad_amax_state=grad_amax_xo)
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [out])
     return (out, *new_amaxs, *saves)
 
-  def feed_forward(self, x:Tensor, ffn_norm:Tensor, w1:Tensor, w2:Tensor, w3:Tensor,
-                   amax_x1=None, amax_w1=None, amax_x2=None, amax_w2=None, amax_x3=None, amax_w3=None):
+  def feed_forward(self, x:Tensor, residual:Tensor, ffn_norm:Tensor, w13:Tensor, w2:Tensor,
+                   amax_x13:Tensor, amax_x2:Tensor, s_13:Tensor, s_2:Tensor,
+                   grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
     new_amaxs, saves = [], []
 
-    x, rrms = rmsnorm(x, self.norm_eps)
-    saves.extend([x, rrms])
-    x = x * ffn_norm
+    x_w13, h, x_normed, rrms, ret = add_norm_quantize_matmul(x, residual, ffn_norm, w13, s_13, self.norm_eps,
+                                                             amax_x=amax_x13)
+    saves.extend([x_normed, rrms])
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [x_w13])
 
-    x_w1, *ret = matmul(x, w1, amax_x=amax_x1, amax_w=amax_w1)
-    new_amaxs.extend(ret[:2])
-    saves.extend(ret[2:] + [x_w1])
-    x_w3, *ret = matmul(x.contiguous_backward(), w3, amax_x=amax_x3, amax_w=amax_w3)
-    new_amaxs.extend(ret[:2])
-    saves.extend(ret[2:] + [x_w3])
-    out, *ret = matmul(x_w1.silu() * x_w3, w2, amax_x=amax_x2, amax_w=amax_w2)
-    new_amaxs.extend(ret[:2])
-    saves.extend(ret[2:] + [out])
-    return (out, *new_amaxs, *saves)
+    out, ret = silu_w13_quantize_matmul(x_w13, w2, s_2, amax_x2=amax_x2, grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
+    new_amaxs.extend(ret[:1])
+    saves.extend(ret[1:] + [out])
+    return (out, h, *new_amaxs, *saves)
 
   @function(precompile=True, precompile_backward=True)
   def run_layer(self, x:Tensor, freqs_cis:Tensor,
                 attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
-                ffn_norm:Tensor, w1:Tensor, w2:Tensor, w3:Tensor,
-                amax_xqkv=None, amax_wqkv=None, amax_xo=None, amax_wo=None,
-                amax_x1=None, amax_w1=None, amax_x2=None, amax_w2=None, amax_x3=None, amax_w3=None):
+                ffn_norm:Tensor, w13:Tensor, w2:Tensor,
+                amax_xqkv:Tensor, amax_xo:Tensor,
+                amax_x13:Tensor, amax_x2:Tensor,
+                s_qkv:Tensor, s_o:Tensor, s_13:Tensor, s_2:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor,
+                grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
     attn, *attn_ret = self.attention(x, freqs_cis, attention_norm, wqkv, wo,
-                                     amax_xqkv=amax_xqkv, amax_wqkv=amax_wqkv, amax_xo=amax_xo, amax_wo=amax_wo)
-    attn_amaxs, attn_saves = attn_ret[:4], attn_ret[4:]
-    h = x + attn
-    ffn, *ffn_ret = self.feed_forward(h, ffn_norm, w1, w2, w3,
-                                      amax_x1=amax_x1, amax_w1=amax_w1, amax_x2=amax_x2, amax_w2=amax_w2, amax_x3=amax_x3, amax_w3=amax_w3)
-    ffn_amaxs, ffn_saves = ffn_ret[:6], ffn_ret[6:]
+                                     amax_xqkv=amax_xqkv, amax_xo=amax_xo, s_qkv=s_qkv, s_o=s_o,
+                                     grad_amax_xqkv=grad_amax_xqkv, grad_amax_xo=grad_amax_xo)
+    attn_amaxs, attn_saves = attn_ret[:2], attn_ret[2:]
+    ffn, h, *ffn_ret = self.feed_forward(x, attn, ffn_norm, w13, w2,
+                                                   amax_x13=amax_x13, amax_x2=amax_x2, s_13=s_13, s_2=s_2,
+                                                   grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
+    ffn_amaxs, ffn_saves = ffn_ret[:2], ffn_ret[2:]
     h = h + ffn
     return (h, *attn_amaxs, *ffn_amaxs, *attn_saves, *ffn_saves)
 
@@ -195,42 +215,40 @@ class FlatTransformer:
     else:
       # flat per-layer weights: axis 0 is n_layers, so shard axes are +1 vs per-layer Transformer
       self.wqkv.shard_(device, axis=1).realize()          # (n_layers, out, dim) shard out
-      self.wo.shard_(device, axis=2).realize()             # (n_layers, dim, in) shard in
-      self.w1.shard_(device, axis=1).realize()             # (n_layers, hidden, dim) shard out
-      self.w2.shard_(device, axis=2).realize()             # (n_layers, dim, hidden) shard in
-      self.w3.shard_(device, axis=1).realize()             # (n_layers, hidden, dim) shard out
+      self.wo.shard_(device, axis=2).realize()            # (n_layers, dim, in) shard in
+      self.w13.shard_(device, axis=1).realize()           # (n_layers, hidden*2, dim) shard out
+      self.w2.shard_(device, axis=2).realize()            # (n_layers, dim, hidden) shard in
       self.attention_norm.shard_(device, axis=None).realize()
       self.ffn_norm.shard_(device, axis=None).realize()
       self.norm.weight.shard_(device, axis=None).realize()
       self.tok_embeddings.weight.shard_(device, axis=0).realize()
       self.output.shard_(device, axis=1).realize()
       self.freqs_cis.shard_(device, axis=None).realize()
-      if FP8:
-        for name in self._fp8_amax:
-          for i in range(len(self._fp8_amax[name])):
-            self._fp8_amax[name][i] = self._fp8_amax[name][i].to(device).contiguous().requires_grad_(False)
+      for amax_dict in (self._fp8_amax, self._fp8_grad_amax):
+        for name in amax_dict:
+          for i in range(len(amax_dict[name])):
+            amax_dict[name][i] = amax_dict[name][i].to(device).contiguous().requires_grad_(False)
+      for name in self._fp8_inv_scale:
+        self._fp8_inv_scale[name] = self._fp8_inv_scale[name].to(device).contiguous().requires_grad_(False)
 
   def __call__(self, tokens:Tensor):
     h = self.tok_embeddings(tokens)
     freqs_cis = self.freqs_cis.cast(h.dtype)[:, :tokens.shape[1], :, :, :]
-    a = self._fp8_amax if FP8 else None
+    a, ga, s = self._fp8_amax, self._fp8_grad_amax, self._fp8_inv_scale
     for i in range(self.n_layers):
-      amax_layer = {"amax_xqkv": a["xqkv"][i], "amax_wqkv": a["wqkv"][i],
-                    "amax_xo": a["xo"][i], "amax_wo": a["wo"][i],
-                    "amax_x1": a["x1"][i], "amax_w1": a["w1"][i],
-                    "amax_x2": a["x2"][i], "amax_w2": a["w2"][i],
-                    "amax_x3": a["x3"][i], "amax_w3": a["w3"][i]} if a else {}
       h, *ret = self.run_layer(h, freqs_cis,
                                self.attention_norm[i], self.wqkv[i], self.wo[i],
-                               self.ffn_norm[i], self.w1[i], self.w2[i], self.w3[i],
-                               **amax_layer)
-      if a:
-        amaxs = ret[:10]
-        amax_names = ["xqkv", "wqkv", "xo", "wo", "x1", "w1", "x3", "w3", "x2", "w2"]
-        for name, new_val in zip(amax_names, amaxs):
-          a[name][i].assign(new_val)
+                               self.ffn_norm[i], self.w13[i], self.w2[i],
+                               amax_xqkv=a["xqkv"][i], amax_xo=a["xo"][i],
+                               amax_x13=a["x13"][i], amax_x2=a["x2"][i],
+                               s_qkv=s["wqkv"][i], s_o=s["wo"][i],
+                               s_13=s["w13"][i], s_2=s["w2"][i],
+                               grad_amax_xqkv=ga["xqkv"][i], grad_amax_xo=ga["xo"][i],
+                               grad_amax_xw13=ga["xw13"][i], grad_amax_xout=ga["xout"][i])
+      for name, new_val in zip(["xqkv", "xo", "x13", "x2"], ret[:5]):
+        a[name][i].assign(new_val)
 
-    logits = matmul(self.norm(h).contiguous().contiguous_backward(), self.output[0], fp8=False)[0].contiguous_backward()
+    logits = matmul(self.norm(h), self.output[0], fp8=False)[0]
     return logits
 
 def _get_pads(uop:UOp) -> list[UOp]:
@@ -239,13 +257,19 @@ def _get_pads(uop:UOp) -> list[UOp]:
 
 def apply_grad(grad_buf:Tensor, new_grad:UOp):
   pads = _get_pads(new_grad)
-  new_grad = new_grad.cast(grad_buf.dtype)
   if len(pads) <= 1:
+    new_grad = new_grad.cast(grad_buf.dtype)
     store = grad_buf.uop.store(grad_buf.uop + new_grad)
     grad_buf.uop = grad_buf.uop.after(store)
     return
   sorted_pads = sorted(pads, key=lambda p: p.marg[0][0] if p.op == Ops.PAD else 0)
-  inners = [Tensor(p.src[0] if p.op == Ops.PAD else p, device=grad_buf.device).cast(grad_buf.dtype) for p in sorted_pads]
+  inners_raw = [Tensor(p.src[0] if p.op == Ops.PAD else p, device=grad_buf.device) for p in sorted_pads]
+  if getenv("FUSED_PAD_GRAD_ACCUM", 0):
+    from extra.llama_kernels.fused_pad_grad_accum import fused_pad_grad_accum, can_fused_pad_grad_accum
+    if can_fused_pad_grad_accum(grad_buf, inners_raw):
+      grad_buf.uop = fused_pad_grad_accum(grad_buf, inners_raw).uop
+      return
+  inners = [t.cast(grad_buf.dtype) for t in inners_raw]
   grad_buf.assign(grad_buf + inners[0].cat(*inners[1:], dim=0))
 
 if __name__ == "__main__":

examples/mlperf/optim.py

@@ -34,7 +34,9 @@ class GradAccClipAdamW(Optimizer):
     else:
       updates, extra = self._step([], grads)
     for i, tt in enumerate(self.params): tt.assign(self._apply_update(tt, updates[i], self.master_params[i] if self.master_params else None))
-    to_realize = extra+self.params+self.buffers+(self.master_params or [])
+    # collect inv_scale tensors attached to fp8 params (set by _apply_update)
+    fp8_inv_scales = [tt._inv_scale for tt in self.params if hasattr(tt, '_inv_scale')]
+    to_realize = extra+self.params+self.buffers+(self.master_params or [])+fp8_inv_scales
 
     Tensor.realize(*to_realize)
     return extra[-1]
@@ -77,4 +79,12 @@ class GradAccClipAdamW(Optimizer):
     new_w = w.detach() - up
     if master is not None: master.assign(new_w)
     if STOCHASTIC_ROUND and t.dtype == dtypes.bfloat16: return stochastic_round_bf16(new_w)
+    if t.dtype in dtypes.fp8s:
+      from examples.mlperf.models.flat_llama import FP8_MAX
+      amax = new_w.float().abs().max(axis=tuple(range(1, new_w.ndim))).detach()  # per-layer amax for (n_layers, out, in)
+      scale = FP8_MAX / (amax + 1e-8)
+      fp8_w = (new_w * scale.reshape(-1, *([1]*(new_w.ndim-1)))).clamp(-FP8_MAX, FP8_MAX).cast(t.dtype)
+      if hasattr(t, '_inv_scale'):
+        t._inv_scale.assign(((amax + 1e-8) / FP8_MAX).cast(t._inv_scale.dtype))
+      return fp8_w
     return new_w.cast(t.dtype)

examples/mlperf/training_submission_v6.0/.gitignore

@@ -0,0 +1 @@
+!*.txt

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_1xMI300X/dev_beam.sh

@@ -1,17 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=1 BS=128 EVAL_BS=128
-
-export CHECK_OOB=0
-
-export BEAM=3 BEAM_UOPS_MAX=4000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-# export BEAM_LOG_SURPASS_MAX=1
-# export BASEDIR="/raid/datasets/wiki"
-
-export RESET_STEP=1
-export BENCHMARK=10 BERT_LAYERS=2 DEBUG=2
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/README.md

@@ -1,69 +0,0 @@
-# 1. Problem
-
-This problem uses BERT for NLP.
-
-## Requirements
-
-Install tinygrad and mlperf-logging (uncomment mlperf from setup.py) from branch mlperf_training_v5.0.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-Also install gdown (for dataset), numpy, tqdm and tensorflow.
-```
-pip install gdown numpy tqdm tensorflow
-```
-
-### 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.
-
-# 2. Directions
-
-## Steps to download and verify data
-
-### 1. Download raw data
-
-```
-BASEDIR="/raid/datasets/wiki" WIKI_TRAIN=1 VERIFY_CHECKSUM=1 python3 extra/datasets/wikipedia_download.py
-```
-
-### 2. Preprocess train and validation data
-
-Note: The number of threads used for preprocessing is limited by available memory. With 128GB of RAM, a maximum of 16 threads is recommended. 
-
-#### Training:
-```
-BASEDIR="/raid/datasets/wiki" NUM_WORKERS=16 python3 extra/datasets/wikipedia.py pre-train all
-```
-
-Generating a specific topic (Between 0 and 499)
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-train 42
-```
-
-#### Validation:
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-eval
-```
-## Running
-
-### tinybox_green
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_green/run_and_time.sh
-```
-
-### tinybox_red
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_red/run_and_time.sh
-```
-### tinybox_8xMI300X
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/run_and_time.sh
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/dev_beam.sh

@@ -1,17 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=3 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-
-export BENCHMARK=10 BERT_LAYERS=2
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/dev_run.sh

@@ -1,20 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-
-# similar to https://github.com/mlcommons/training_results_v3.1/blob/d06288b2bd675a9d88e0e6181f5bb5626b71ec19/Quanta_Cloud_Technology/results/D54U-3U/bert/result_1.txt#L54
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-export TRAIN_STEPS=3900
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=3 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-
-export WANDB=1 PARALLEL=0
-
-RUNMLPERF=1 python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/run_and_time.sh

@@ -1,31 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export SUBMISSION_PLATFORM="tinybox_8xMI300X"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-
-# similar to https://github.com/mlcommons/training_results_v3.1/blob/d06288b2bd675a9d88e0e6181f5bb5626b71ec19/Quanta_Cloud_Technology/results/D54U-3U/bert/result_1.txt#L54
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-export TRAIN_STEPS=3900
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=3 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="bert_8xMI300x_${DATETIME}_${SEED}.log"
-
-BENCHMARK=10 INITMLPERF=1 BERT_LAYERS=2 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI350X/dev_run_fa.sh

@@ -1,20 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD DEBUG=0 JIT=1 FLASH_ATTENTION=1
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-
-# similar to https://github.com/mlcommons/training_results_v3.1/blob/d06288b2bd675a9d88e0e6181f5bb5626b71ec19/Quanta_Cloud_Technology/results/D54U-3U/bert/result_1.txt#L54
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-export TRAIN_STEPS=3900
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000
-
-export BEAM=0 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-
-export WANDB=1 PARALLEL=0
-
-RUNMLPERF=1 python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI350X/dev_run_fp8.sh

@@ -1,24 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-
-# similar to https://github.com/mlcommons/training_results_v3.1/blob/d06288b2bd675a9d88e0e6181f5bb5626b71ec19/Quanta_Cloud_Technology/results/D54U-3U/bert/result_1.txt#L54
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-export TRAIN_STEPS=3900
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000
-
-export BEAM=3 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-export BEAM_TIMEOUT_SEC=15
-export FP8_TRAIN=1
-# search
-IGNORE_BEAM_CACHE=1 BENCHMARK=10 BERT_LAYERS=2 RUNMLPERF=0 python3 examples/mlperf/model_train.py
-
-export WANDB=1 PARALLEL=0
-
-RUNMLPERF=1 python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI350X/run_and_time.sh

@@ -1,31 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export SUBMISSION_PLATFORM="tinybox_8xMI350X"
-export DEFAULT_FLOAT="HALF" GPUS=8 BS=1024 EVAL_BS=1024
-
-# similar to https://github.com/mlcommons/training_results_v3.1/blob/d06288b2bd675a9d88e0e6181f5bb5626b71ec19/Quanta_Cloud_Technology/results/D54U-3U/bert/result_1.txt#L54
-export OPT_BASE_LEARNING_RATE=0.0011 OPT_LAMB_BETA_1=0.60466 OPT_LAMB_BETA_2=0.85437 DECAY=0.1
-export TRAIN_STEPS=3900
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000
-
-export BEAM=3 BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1 FREE_INTERMEDIATE=0
-export BASEDIR="/raid/datasets/wiki"
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="bert_8xMI350x_${DATETIME}_${SEED}.log"
-
-BENCHMARK=10 INITMLPERF=1 BERT_LAYERS=2 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_green/README.md

@@ -1,69 +0,0 @@
-# 1. Problem
-
-This problem uses BERT for NLP.
-
-## Requirements
-
-Install tinygrad and mlperf-logging (uncomment mlperf from setup.py) from branch mlperf_training_v5.0.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-Also install gdown (for dataset), numpy, tqdm and tensorflow.
-```
-pip install gdown numpy tqdm tensorflow
-```
-
-### 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.
-
-# 2. Directions
-
-## Steps to download and verify data
-
-### 1. Download raw data
-
-```
-BASEDIR="/raid/datasets/wiki" WIKI_TRAIN=1 VERIFY_CHECKSUM=1 python3 extra/datasets/wikipedia_download.py
-```
-
-### 2. Preprocess train and validation data
-
-Note: The number of threads used for preprocessing is limited by available memory. With 128GB of RAM, a maximum of 16 threads is recommended. 
-
-#### Training:
-```
-BASEDIR="/raid/datasets/wiki" NUM_WORKERS=16 python3 extra/datasets/wikipedia.py pre-train all
-```
-
-Generating a specific topic (Between 0 and 499)
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-train 42
-```
-
-#### Validation:
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-eval
-```
-## Running
-
-### tinybox_green
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_green/run_and_time.sh
-```
-
-### tinybox_red
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_red/run_and_time.sh
-```
-### tinybox_8xMI300X
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/run_and_time.sh
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_green/dev_beam.sh

@@ -1,17 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=72 EVAL_BS=72
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=8 BEAM_UOPS_MAX=10000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BEAM_LOG_SURPASS_MAX=1
-export BASEDIR="/raid/datasets/wiki"
-
-export BENCHMARK=10 BERT_LAYERS=2 DEBUG=2
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_green/dev_run.sh

@@ -1,16 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=72 EVAL_BS=72
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=8 BEAM_UOPS_MAX=10000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BASEDIR="/raid/datasets/wiki"
-
-export WANDB=1 PARALLEL=0
-
-RUNMLPERF=1 python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_green/run_and_time.sh

@@ -1,28 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=NV
-export MODEL="bert"
-export SUBMISSION_PLATFORM="tinybox_green"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=72 EVAL_BS=72
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=8 BEAM_UOPS_MAX=10000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BASEDIR="/raid/datasets/wiki"
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="bert_green_${DATETIME}_${SEED}.log"
-
-# init
-BENCHMARK=10 INITMLPERF=1 BERT_LAYERS=2 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_red/README.md

@@ -1,69 +0,0 @@
-# 1. Problem
-
-This problem uses BERT for NLP.
-
-## Requirements
-
-Install tinygrad and mlperf-logging (uncomment mlperf from setup.py) from branch mlperf_training_v5.0.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-Also install gdown (for dataset), numpy, tqdm and tensorflow.
-```
-pip install gdown numpy tqdm tensorflow
-```
-
-### 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.
-
-# 2. Directions
-
-## Steps to download and verify data
-
-### 1. Download raw data
-
-```
-BASEDIR="/raid/datasets/wiki" WIKI_TRAIN=1 VERIFY_CHECKSUM=1 python3 extra/datasets/wikipedia_download.py
-```
-
-### 2. Preprocess train and validation data
-
-Note: The number of threads used for preprocessing is limited by available memory. With 128GB of RAM, a maximum of 16 threads is recommended. 
-
-#### Training:
-```
-BASEDIR="/raid/datasets/wiki" NUM_WORKERS=16 python3 extra/datasets/wikipedia.py pre-train all
-```
-
-Generating a specific topic (Between 0 and 499)
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-train 42
-```
-
-#### Validation:
-```
-BASEDIR="/raid/datasets/wiki" python3 extra/datasets/wikipedia.py pre-eval
-```
-## Running
-
-### tinybox_green
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_green/run_and_time.sh
-```
-
-### tinybox_red
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_red/run_and_time.sh
-```
-### tinybox_8xMI300X
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/bert/implementations/tinybox_8xMI300X/run_and_time.sh
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_red/dev_beam.sh

@@ -1,18 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=96 EVAL_BS=96
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=5 BEAM_UOPS_MAX=8000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BEAM_LOG_SURPASS_MAX=1
-export BASEDIR="/raid/datasets/wiki"
-
-export RESET_STEP=1
-export BENCHMARK=10 BERT_LAYERS=2 DEBUG=2
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_red/dev_run.sh

@@ -1,16 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=96 EVAL_BS=96
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=5 BEAM_UOPS_MAX=8000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BASEDIR="/raid/datasets/wiki"
-
-export WANDB=1 PARALLEL=0
-
-RUNMLPERF=1 python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/bert/implementations/tinybox_red/run_and_time.sh

@@ -1,31 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="bert"
-export SUBMISSION_PLATFORM="tinybox_red"
-export DEFAULT_FLOAT="HALF" SUM_DTYPE="HALF" GPUS=6 BS=96 EVAL_BS=96
-
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=500000
-
-export BEAM=5 BEAM_UOPS_MAX=8000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5
-export IGNORE_JIT_FIRST_BEAM=1
-export BASEDIR="/raid/datasets/wiki"
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=${LOGMLPERF:-1}
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="bert_red_${DATETIME}_${SEED}.log"
-
-export HCQDEV_WAIT_TIMEOUT_MS=100000  # prevents hang?
-
-# init
-sleep 5 && sudo rmmod amdgpu || true
-BENCHMARK=10 INITMLPERF=1 BERT_LAYERS=2 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/README.md

@@ -0,0 +1,28 @@
+# 1. Problem
+
+small llm pretraining: llama 3.1 8b on c4.
+
+## Requirements
+
+Install tinygrad and mlperf-logging (uncomment mlperf from setup.py) from branch mlperf_training_v6.0.
+```
+git clone https://github.com/tinygrad/tinygrad.git
+python3 -m pip install -e ".[mlperf]"
+```
+
+# 2. Directions
+
+## Steps to download and verify data
+
+### 1. Download raw data
+
+follow mlperf steps to download the preprocessed c4 dataset.
+
+## Running
+
+### tinybox_8xMI350X
+
+#### Steps to run benchmark
+```
+examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/run_and_time.sh
+```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/dev_beam.sh

@@ -9,15 +9,21 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-2}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
+export LATE_ALLREDUCE=${LATE_ALLREDUCE:-0}
 export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
 export WQKV=${WQKV:-1}
 export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
 export FP8=${FP8:-1}
 export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-16} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
+export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
 export GBS=$((BS * GRADIENT_ACC_STEPS))
 
 export MODEL="llama3"
@@ -36,7 +42,7 @@ export DATA_SEED=${DATA_SEED:-5760}
 export JITBEAM=${JITBEAM:-3}
 export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
 
-export FAKEDATA=1 BENCHMARK=${BENCHMARK:-10}
+export FAKEDATA=${FAKEDATA:-1} BENCHMARK=${BENCHMARK:-10}
 if [ -z "$FULL_LAYERS" ]; then
   export LLAMA_LAYERS=2
 fi

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/dev_run.sh

@@ -9,15 +9,21 @@ export DEVICE_IN_FUNCTION_BUG=1
 export DEBUG=${DEBUG:-0}
 export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
 export ALL2ALL=${ALL2ALL:-1}
+export LATE_ALLREDUCE=${LATE_ALLREDUCE:-0}
 export USE_ATOMICS=${USE_ATOMICS:-1}
 export ASM_GEMM=${ASM_GEMM:-1}
 export WQKV=${WQKV:-1}
 export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
 export FP8=${FP8:-1}
 export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
+export FAST_CE=${FAST_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-16} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
+export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
 export GBS=$((BS * GRADIENT_ACC_STEPS))
 
 export MODEL="llama3"

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/profile.sh

@@ -3,4 +3,4 @@ export BENCHMARK=5
 export EVAL_BS=0
 VIZ=${VIZ:--1} FULL_LAYERS=1 DEBUG=0 examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_beam.sh
 SRC="AMD"; [[ $DEV == NULL* ]] && SRC="NULL"
-extra/viz/cli.py --profile -s "$SRC"
+python -m tinygrad.viz.cli -s "$SRC" -t

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/run_and_time.sh

@@ -10,15 +10,21 @@ export DEVICE_IN_FUNCTION_BUG=1
 
 export HK_FLASH_ATTENTION=1
 export ALL2ALL=1
+export LATE_ALLREDUCE=0
 export USE_ATOMICS=1
 export ASM_GEMM=1
 export WQKV=1
 export MASTER_WEIGHTS=1
 export FP8=1
 export ALLREDUCE_CAST=1
+export FAST_CE=1
+export FUSED_INPUT_QUANTIZE=1
+export FUSED_ADD_NORM_MUL_QUANTIZE=1
+export FUSED_SILU_W13=1
+export FUSED_PAD_GRAD_ACCUM=1
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=8 MP=1 BS=16 EVAL_BS=16 GRADIENT_ACC_STEPS=2
+export DP=8 MP=1 BS=16 EVAL_BS=8 GRADIENT_ACC_STEPS=2
 export GBS=$((BS * GRADIENT_ACC_STEPS))
 
 export MODEL="llama3"

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama405b/implementations/tinybox_8xMI350X/dev_beam.sh

@@ -1,38 +0,0 @@
-#!/usr/bin/env bash
-
-export PYTHONPATH="."
-export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
-export DEVICE_IN_FUNCTION_BUG=1
-
-export DEBUG=${DEBUG:-2}
-export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
-export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
-export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-1}
-export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}
-
-export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-1} MP=${MP:-8}
-export BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
-
-export MODEL="llama3"
-export BASEDIR="/raid/datasets/c4/"
-export LLAMA3_SIZE=${LLAMA3_SIZE:-"405B"}
-export SEQLEN=${SEQLEN:-8192}
-
-export SEED=${SEED:-5760}
-export DATA_SEED=${DATA_SEED:-5760}
-
-export JITBEAM=${JITBEAM:-3}
-export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
-
-export FAKEDATA=1 BENCHMARK=10
-if [ -z "$FULL_LAYERS" ]; then
-  export LLAMA_LAYERS=2
-fi
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama405b/implementations/tinybox_8xMI350X/dev_run.sh

@@ -1,32 +0,0 @@
-#!/usr/bin/env bash
-
-export PYTHONPATH="."
-export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
-
-export DEBUG=${DEBUG:-0}
-export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
-export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
-export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-1}
-export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}
-
-export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-1} MP=${MP:-8}
-export BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-1152}
-
-export MODEL="llama3"
-export BASEDIR="/raid/datasets/c4/"
-export LLAMA3_SIZE=${LLAMA3_SIZE:-"405B"}
-export SEQLEN=${SEQLEN:-8192}
-
-export SEED=${SEED:-$RANDOM}
-export DATA_SEED=${DATA_SEED:-5760}
-
-export JITBEAM=${JITBEAM:-3}
-export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_beam_mp.sh

@@ -1,43 +0,0 @@
-#!/usr/bin/env bash
-
-export PYTHONPATH="."
-export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
-export DEVICE_IN_FUNCTION_BUG=1
-
-export DEBUG=${DEBUG:-2}
-export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
-export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
-export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-1}
-export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}
-
-export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-1} MP=${MP:-8} BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}
-export GBS=$((BS * GRADIENT_ACC_STEPS))
-
-export MODEL="llama3"
-export BASEDIR="/raid/datasets/c4-8b/"
-export SMALL=1
-export LLAMA3_SIZE=${LLAMA3_SIZE:-"8B"}
-export EVAL_TARGET=3.3 EVAL_FREQ=12288
-export LR="1e-3" END_LR="1e-4" WARMUP_SAMPLES=4096 MAX_STEPS=1200000
-export WARMUP_STEPS=$((WARMUP_SAMPLES / GBS))
-export SAMPLES=$((MAX_STEPS * GBS))
-export SEQLEN=${SEQLEN:-8192}
-
-export SEED=${SEED:-5760}
-export DATA_SEED=${DATA_SEED:-5760}
-
-export JITBEAM=${JITBEAM:-3}
-export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
-
-export FAKEDATA=1 BENCHMARK=10
-if [ -z "$FULL_LAYERS" ]; then
-  export LLAMA_LAYERS=2
-fi
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_run_mp.sh

@@ -1,38 +0,0 @@
-#!/usr/bin/env bash
-
-export PYTHONPATH="."
-export DEV=${DEV:-AMD}
-export EMULATE="AMD_CDNA4"
-export CHECK_OOB=0
-export REWRITE_STACK_LIMIT=5000000 HCQDEV_WAIT_TIMEOUT_MS=240000
-export DEVICE_IN_FUNCTION_BUG=1
-
-export DEBUG=${DEBUG:-0}
-export HK_FLASH_ATTENTION=${HK_FLASH_ATTENTION:-1}
-export ALL2ALL=${ALL2ALL:-1}
-export USE_ATOMICS=${USE_ATOMICS:-0}
-export ASM_GEMM=${ASM_GEMM:-1}
-export WQKV=${WQKV:-1}
-export OFFLOAD_OPTIM=${OFFLOAD_OPTIM:-1}
-
-export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
-export DP=${DP:-1} MP=${MP:-8} BS=${BS:-1} EVAL_BS=${EVAL_BS:-1} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-32}
-export GBS=$((BS * GRADIENT_ACC_STEPS))
-
-export MODEL="llama3"
-export BASEDIR="/raid/datasets/c4-8b/"
-export SMALL=1
-export LLAMA3_SIZE=${LLAMA3_SIZE:-"8B"}
-export EVAL_TARGET=3.3 EVAL_FREQ=12288
-export LR="1e-3" END_LR="1e-4" WARMUP_SAMPLES=4096 MAX_STEPS=1200000
-export WARMUP_STEPS=$((WARMUP_SAMPLES / GBS))
-export SAMPLES=$((MAX_STEPS * GBS))
-export SEQLEN=${SEQLEN:-8192}
-
-export SEED=${SEED:-$RANDOM}
-export DATA_SEED=${DATA_SEED:-5760}
-
-export JITBEAM=${JITBEAM:-3}
-export BEAM_UOPS_MAX=6000 BEAM_UPCAST_MAX=256 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=1
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_green/README.md

@@ -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
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_green/dev_beam.sh

@@ -1,13 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_green/dev_run.sh

@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_green/run_and_time.sh

@@ -1,25 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=NV
-export MODEL="resnet"
-export SUBMISSION_PLATFORM="tinybox_green"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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=${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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/README.md

@@ -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
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/dev_beam.sh

@@ -1,13 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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=${DEBUG:-2}
-
-python3 examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/dev_run.sh

@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="resnet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/run_and_time.sh

@@ -1,26 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="resnet"
-export SUBMISSION_PLATFORM="tinybox_red"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=1536 EVAL_BS=192
-
-export 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=${LOGMLPERF:-1}
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="resnet_red_${DATETIME}_${SEED}.log"
-
-# init
-sleep 5 && sudo rmmod amdgpu || true
-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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/resnet/implementations/tinybox_red/setup.sh

@@ -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

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_green/README.md

@@ -1,38 +0,0 @@
-# 1. Problem
-
-This problem uses RetinaNet for SSD.
-
-## Requirements
-
-Install tinygrad and mlperf-logging (uncomment mlperf from setup.py) from branch mlperf_training_v5.0.
-```
-git clone https://github.com/tinygrad/tinygrad.git
-python3 -m pip install -e ".[mlperf]"
-```
-
-Also install the following dependencies:
-```
-pip install tqdm numpy pycocotools boto3 pandas torch torchvision
-```
-
-### 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.
-
-# 2. Directions
-
-## Steps to download data
-
-Run the following:
-```
-BASEDIR=/raid/datasets/openimages python3 extra/datasets/openimages.py
-```
-
-## Running
-
-### tinybox_green
-
-#### Steps to run benchmark
-```
-examples/mlperf/training_submission_v5.0/tinycorp/benchmarks/retinanet/implementations/tinybox_green/run_and_time.sh
-```

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_green/dev_beam.sh

@@ -1,14 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="retinanet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=96 EVAL_BS=96
-export BASEDIR="/raid/datasets/openimages"
-
-# export RESET_STEP=0
-
-export TRAIN_BEAM=2 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export BENCHMARK=5 DEBUG=2
-
-python examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_green/dev_run.sh

@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=NV
-export MODEL="retinanet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=96 EVAL_BS=96
-export BASEDIR="/raid/datasets/openimages"
-
-# export RESET_STEP=0
-
-export TRAIN_BEAM=2 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export WANDB=1 PARALLEL=0
-export RUNMLPERF=1
-
-python examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_green/run_and_time.sh

@@ -1,25 +0,0 @@
-#!/bin/bash
-set -e  # Exit on any error
-set -o pipefail  # Make pipeline fail if any command fails
-
-export PYTHONPATH="." DEV=NV
-export MODEL="retinanet"
-export SUBMISSION_PLATFORM="tinybox_green"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=96 EVAL_BS=96
-
-export TRAIN_BEAM=2 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-export IGNORE_JIT_FIRST_BEAM=1
-export BASEDIR="/raid/datasets/openimages"
-
-# pip install -e ".[mlperf]"
-export LOGMLPERF=1
-
-export SEED=$RANDOM
-DATETIME=$(date "+%m%d%H%M")
-LOGFILE="retinanet_green_${DATETIME}_${SEED}.log"
-
-# init
-BENCHMARK=10 INITMLPERF=1 python3 examples/mlperf/model_train.py | tee $LOGFILE
-
-# run
-PARALLEL=0 RUNMLPERF=1 python3 examples/mlperf/model_train.py | tee -a $LOGFILE

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_red/dev_beam.sh

@@ -1,14 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="retinanet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=96 EVAL_BS=96
-export BASEDIR="/raid/datasets/openimages"
-
-# export RESET_STEP=0
-
-export TRAIN_BEAM=2 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export BENCHMARK=5 DEBUG=2
-
-python examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/retinanet/implementations/tinybox_red/dev_run.sh

@@ -1,15 +0,0 @@
-#!/bin/bash
-
-export PYTHONPATH="." DEV=AMD
-export MODEL="retinanet"
-export DEFAULT_FLOAT="HALF" GPUS=6 BS=96 EVAL_BS=96
-export BASEDIR="/raid/datasets/openimages"
-
-# export RESET_STEP=0
-
-export TRAIN_BEAM=2 IGNORE_JIT_FIRST_BEAM=1 BEAM_UOPS_MAX=1500 BEAM_UPCAST_MAX=64 BEAM_LOCAL_MAX=1024 BEAM_MIN_PROGRESS=5 BEAM_PADTO=0
-
-export WANDB=1 PARALLEL=0
-export RUNMLPERF=1
-
-python examples/mlperf/model_train.py

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_0.txt

@@ -0,0 +1,106 @@
+:::MLLOG {"namespace": "", "time_ms": 1778207373785, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373789, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373790, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373790, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373790, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373791, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778207373791, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778207734506, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747904, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "seed", "value": 25580, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747908, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778207747909, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778208080716, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778208080717, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778208901302, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778208901303, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778208952059, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.705078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778208952060, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778208952060, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778209608282, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778209608282, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778209637796, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.552001953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778209637796, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778209637797, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778210294879, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778210294879, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778210324584, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.1011962890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778210324584, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778210324585, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778210980564, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778210980565, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778211010225, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.8807373046875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778211010225, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778211010226, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778211667184, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778211667185, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778211696784, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.7498779296875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778211696785, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778211696786, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778212356059, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778212356060, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778212385775, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.65478515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778212385776, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778212385776, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778213044774, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778213044775, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778213074311, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5731201171875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778213074312, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778213074313, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778213732225, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778213732225, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778213761806, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5137939453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778213761806, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778213761807, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778214419768, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778214419769, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778214449443, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.46630859375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778214449444, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778214449445, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778215112018, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778215112019, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778215141586, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.428955078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778215141586, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778215141587, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778215794970, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778215794970, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778215824346, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.390869140625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778215824346, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778215824347, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778216475810, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778216475810, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778216505269, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.361328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778216505269, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778216505270, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778217157389, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778217157390, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778217186831, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.346923828125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778217186832, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778217186832, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778217846265, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778217846266, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778217876013, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3133544921875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778217876014, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778217876014, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778218532377, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778218532378, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778218561863, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2989501953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778218561863, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778218561864, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778218561864, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_1.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778218577779, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778218577783, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778218577784, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778218577784, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778218577784, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778218578371, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778218578371, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778218957180, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971058, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "seed", "value": 356, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971063, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778218971064, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778219289653, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778219289654, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778220097041, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778220097042, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778220141757, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.743896484375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778220141758, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778220141758, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778220795772, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778220795773, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778220825439, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.58349609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778220825440, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778220825440, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778221480609, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778221480610, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778221510284, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.1131591796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778221510285, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778221510286, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778222164664, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778222164665, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778222194290, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.8935546875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778222194291, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778222194291, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778222848846, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778222848847, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778222878557, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.7567138671875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778222878558, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778222878558, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778223532447, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778223532447, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778223562036, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.658203125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778223562037, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778223562037, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778224215343, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778224215344, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778224244924, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5860595703125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778224244925, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778224244925, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778224898378, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778224898379, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778224928021, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.51708984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778224928021, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778224928022, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778225581424, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778225581425, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778225611002, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.471923828125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778225611003, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778225611003, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778226265043, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778226265044, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778226294659, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.43701171875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778226294660, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778226294661, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778226949577, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778226949577, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778226979238, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5406494140625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778226979239, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778226979239, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778227635352, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778227635352, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778227664978, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3836669921875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778227664978, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778227664979, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778228323150, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778228323151, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778228352865, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.355712890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778228352865, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778228352866, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778229010307, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778229010307, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778229040142, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3319091796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778229040143, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778229040143, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778229696378, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778229696379, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778229726195, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.30615234375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778229726195, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778229726196, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778230383239, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778230383240, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778230412831, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.29052734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778230412832, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778230412832, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778230412833, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_2.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778230427283, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427287, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427287, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427287, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427287, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427939, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778230427939, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778230779581, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792886, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792890, "event_type": "POINT_IN_TIME", "key": "seed", "value": 2774, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792891, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778230792892, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778231115792, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778231115793, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778232030906, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778232030907, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778232075494, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.812255859375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778232075494, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778232075495, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778232729579, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778232729580, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778232759140, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.582275390625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778232759141, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778232759142, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778233413630, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778233413631, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778233443219, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.11767578125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778233443220, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778233443220, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778234097427, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778234097428, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778234127034, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.9005126953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778234127034, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778234127035, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778234780955, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778234780956, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778234810558, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.7586669921875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778234810558, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778234810559, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778235463904, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778235463905, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778235493473, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.657958984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778235493474, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778235493475, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778236147005, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778236147005, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778236176551, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.585693359375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778236176552, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778236176552, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778236830530, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778236830530, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778236860107, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.521484375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778236860108, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778236860108, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778237514002, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778237514003, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778237543592, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4742431640625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778237543592, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778237543593, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778238197935, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778238197936, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778238227501, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.428955078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778238227502, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778238227503, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778238882036, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778238882037, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778238911645, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4019775390625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778238911645, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778238911646, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778239565129, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778239565130, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778239594721, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.37890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778239594722, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778239594722, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778240248763, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778240248764, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778240278335, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3448486328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778240278336, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778240278337, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778240933651, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778240933651, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778240963429, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.325439453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778240963430, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778240963431, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778241626264, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778241626265, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778241656303, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3072509765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778241656304, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778241656304, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778242315322, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778242315323, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778242345178, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2781982421875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778242345178, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778242345179, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778242345179, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_3.txt

@@ -0,0 +1,106 @@
+:::MLLOG {"namespace": "", "time_ms": 1778242359541, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778242359545, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778242359545, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778242359545, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778242359545, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778242360117, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778242360118, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778242702158, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715949, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715953, "event_type": "POINT_IN_TIME", "key": "seed", "value": 1261, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715953, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715953, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715954, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715955, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715955, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715955, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778242715955, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778243033805, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778243033806, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778243851371, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778243851372, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778243896651, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.7802734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778243896652, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778243896652, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778244555628, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778244555629, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778244585531, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.574951171875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778244585532, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778244585533, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778245246511, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778245246512, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778245276502, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.1171875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778245276503, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778245276503, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778245937187, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778245937187, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778245967058, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.8995361328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778245967059, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778245967059, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778246626117, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778246626117, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778246656019, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.762451171875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778246656019, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778246656020, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778247315255, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778247315256, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778247345128, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6572265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778247345128, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778247345129, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778248003582, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778248003582, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778248033442, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.58740234375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778248033443, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778248033443, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778248692764, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778248692764, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778248722726, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5286865234375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778248722727, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778248722727, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778249383186, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778249383186, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778249413099, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.475830078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778249413099, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778249413100, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778250072852, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778250072852, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778250102740, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4278564453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778250102741, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778250102741, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778250762230, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778250762230, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778250792198, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.400146484375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778250792199, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778250792199, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778251455492, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778251455492, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778251485544, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3818359375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778251485545, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778251485545, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778252146772, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778252146772, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778252176776, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.345458984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778252176776, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778252176777, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778252836585, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778252836586, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778252866442, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.322265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778252866443, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778252866443, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778253526422, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778253526422, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778253556343, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.299072265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778253556343, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778253556344, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778253556344, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_4.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778253570454, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778253570459, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778253570459, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778253570459, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778253570459, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778253571045, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778253571045, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778253944036, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957691, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "seed", "value": 14711, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957695, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778253957696, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778254276545, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778254276546, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778255100535, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778255100536, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778255143977, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.77978515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778255143977, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778255143978, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778255806844, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778255806845, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778255836518, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.578857421875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778255836519, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778255836520, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778256495933, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778256495933, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778256525443, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.1239013671875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778256525443, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778256525444, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778257180826, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778257180827, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778257210282, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.906494140625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778257210283, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778257210283, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778257866434, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778257866435, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778257895945, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.75244140625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778257895945, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778257895946, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778258550818, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778258550819, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778258580369, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6553955078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778258580369, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778258580370, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778259234200, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778259234201, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778259263770, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5762939453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778259263771, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778259263772, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778259917494, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778259917495, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778259947011, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.52197265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778259947012, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778259947013, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778260600453, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778260600454, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778260629950, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778260629951, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778260629951, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778261285126, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778261285127, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778261314809, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4378662109375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778261314810, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778261314810, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778261971632, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778261971632, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778262001260, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3968505859375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778262001261, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778262001261, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778262657393, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778262657394, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778262686962, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.365966796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778262686962, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778262686963, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778263342665, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778263342666, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778263372176, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3365478515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778263372176, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778263372177, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778264027427, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778264027428, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778264056993, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3363037109375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778264056993, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778264056994, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778264710992, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778264710993, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778264740486, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3016357421875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778264740486, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778264740487, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778265396989, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778265396989, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778265426521, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2861328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778265426522, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778265426522, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778265426522, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_5.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778265440911, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778265440915, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778265440915, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778265440916, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778265440916, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778265441493, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778265441493, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778265779467, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792765, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "seed", "value": 27754, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792769, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778265792770, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778266108942, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778266108943, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778266913943, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778266913944, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778266957471, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.74072265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778266957472, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778266957472, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778267616663, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778267616663, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778267648052, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.612060546875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778267648053, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778267648053, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778268306168, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778268306168, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778268335863, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.16552734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778268335864, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778268335864, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778268998030, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778268998030, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778269027991, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.915283203125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778269027992, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778269027992, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778269689514, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778269689515, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778269719312, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.7637939453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778269719313, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778269719313, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778270378319, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778270378320, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778270408037, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6695556640625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778270408038, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778270408038, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778271066429, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778271066430, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778271096134, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.583251953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778271096135, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778271096135, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778271754376, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778271754377, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778271784142, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.525146484375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778271784142, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778271784143, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778272442458, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778272442459, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778272472257, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4774169921875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778272472257, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778272472258, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778273129575, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778273129576, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778273159231, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.443359375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778273159231, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778273159232, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778273816098, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778273816099, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778273845769, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4072265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778273845770, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778273845770, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778274505683, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778274505684, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778274535540, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3677978515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778274535541, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778274535541, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778275195662, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778275195662, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778275225396, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4146728515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778275225397, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778275225397, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778275884245, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778275884246, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778275913924, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3697509765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778275913925, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778275913925, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778276570930, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778276570931, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778276600619, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.321533203125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778276600620, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778276600620, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778277262406, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778277262407, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778277292466, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.287353515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778277292467, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778277292467, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778277292468, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_6.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778277306868, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778277306872, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778277306872, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778277306873, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778277306873, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778277307428, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778277307429, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778277671564, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685153, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685157, "event_type": "POINT_IN_TIME", "key": "seed", "value": 17816, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685157, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685157, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685158, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685159, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685159, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778277685159, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778278007248, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778278007260, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778278810368, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778278810369, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778278855284, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.768798828125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778278855285, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778278855285, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778279519460, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778279519461, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778279549391, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.568603515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778279549392, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778279549392, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778280214562, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778280214563, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778280244495, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.151123046875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778280244496, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778280244496, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778280909906, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778280909906, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778280939913, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.9197998046875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778280939913, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778280939914, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778281607749, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778281607750, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778281637814, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.7734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778281637815, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778281637815, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778282306223, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778282306224, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778282336322, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.673583984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778282336323, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778282336323, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778283007699, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778283007700, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778283037808, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6011962890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778283037808, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778283037809, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778283706598, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778283706598, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778283736748, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.526123046875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778283736748, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778283736749, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778284408590, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778284408590, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778284438316, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.475341796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778284438317, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778284438317, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778285098897, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778285098898, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778285128703, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.432861328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778285128703, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778285128704, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778285786660, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778285786660, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778285816222, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4031982421875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778285816222, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778285816223, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778286473781, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778286473782, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778286503417, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3638916015625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778286503418, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778286503418, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778287160556, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778287160556, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778287190213, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.341796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778287190214, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778287190215, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778287846424, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778287846424, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778287876044, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.32177734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778287876045, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778287876046, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778288531947, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778288531947, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778288561549, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5465087890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778288561550, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778288561550, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778289220442, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778289220442, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778289250127, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2855224609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778289250128, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778289250128, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778289250129, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_7.txt

@@ -0,0 +1,106 @@
+:::MLLOG {"namespace": "", "time_ms": 1778289264340, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264344, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264344, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264344, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264344, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264911, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778289264912, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778289599730, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613197, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613200, "event_type": "POINT_IN_TIME", "key": "seed", "value": 16781, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613201, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778289613202, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778289929875, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778289929878, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778290756967, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778290756968, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778290801735, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.758544921875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778290801736, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778290801736, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778291460896, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778291460896, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778291490685, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.683349609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778291490685, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778291490686, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778292152773, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778292152774, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778292182518, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.1280517578125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778292182519, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778292182519, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778292842100, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778292842101, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778292871768, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.90185546875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778292871769, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778292871769, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778293529314, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778293529315, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778293559042, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.757080078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778293559043, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778293559043, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778294218188, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778294218189, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778294247880, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6575927734375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778294247880, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778294247881, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778294908017, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778294908018, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778294937688, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.586181640625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778294937689, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778294937690, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778295595710, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778295595710, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778295625392, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5230712890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778295625393, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778295625394, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778296283795, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778296283795, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778296313518, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.467529296875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778296313519, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778296313519, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778296973892, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778296973893, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778297003579, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4351806640625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778297003580, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778297003580, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778297661577, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778297661578, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778297691130, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.406982421875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778297691130, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778297691131, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778298348217, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778298348218, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778298377837, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3848876953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778298377837, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778298377838, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778299035939, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778299035940, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778299065575, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3480224609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778299065576, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778299065576, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778299724382, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778299724383, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778299754023, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3209228515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778299754023, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778299754024, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778300412415, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778300412415, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778300442058, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2950439453125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778300442059, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778300442060, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778300442060, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_8.txt

@@ -0,0 +1,111 @@
+:::MLLOG {"namespace": "", "time_ms": 1778300456451, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778300456455, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778300456455, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778300456455, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778300456455, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778300457011, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778300457012, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778300803665, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817390, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "seed", "value": 4729, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817395, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778300817396, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778301145773, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778301145774, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778301985088, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778301985089, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778302030319, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.865966796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778302030319, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778302030320, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778302687526, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778302687527, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778302717259, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.615966796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778302717260, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778302717261, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778303376036, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778303376037, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778303406044, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.154296875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778303406045, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778303406045, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778304071224, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778304071225, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778304101168, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.9095458984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778304101169, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778304101170, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778304762172, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778304762173, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778304792161, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.775634765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778304792162, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778304792162, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778305452836, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778305452836, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778305482708, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.676513671875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778305482708, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778305482709, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778306140246, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778306140246, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778306169947, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5947265625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778306169947, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778306169948, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778306828284, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778306828285, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778306858077, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5255126953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778306858077, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778306858078, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778307519609, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778307519610, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778307549531, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4757080078125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778307549532, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778307549532, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778308208151, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778308208152, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778308237856, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4312744140625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778308237857, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778308237857, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778308896397, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778308896398, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778308926271, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.402099609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778308926271, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778308926272, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778309586346, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778309586347, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778309616134, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.37060546875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778309616134, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778309616135, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778310273337, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778310273338, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778310303090, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3968505859375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778310303091, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778310303092, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778310958883, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778310958883, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778310988541, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3284912109375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778310988542, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778310988542, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778311645004, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778311645004, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778311674742, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.302001953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778311674743, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778311674744, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778312331845, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778312331846, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778312361570, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2777099609375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778312361571, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778312361571, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 196608}}
+:::MLLOG {"namespace": "", "time_ms": 1778312361572, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/results/tinybox_8xMI350X/llama31_8b/result_9.txt

@@ -0,0 +1,106 @@
+:::MLLOG {"namespace": "", "time_ms": 1778312377935, "event_type": "POINT_IN_TIME", "key": "submission_org", "value": "tinycorp", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1328}}
+:::MLLOG {"namespace": "", "time_ms": 1778312377940, "event_type": "POINT_IN_TIME", "key": "submission_platform", "value": "tinybox", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1329}}
+:::MLLOG {"namespace": "", "time_ms": 1778312377940, "event_type": "POINT_IN_TIME", "key": "submission_division", "value": "closed", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1330}}
+:::MLLOG {"namespace": "", "time_ms": 1778312377940, "event_type": "POINT_IN_TIME", "key": "submission_status", "value": "onprem", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1331}}
+:::MLLOG {"namespace": "", "time_ms": 1778312377940, "event_type": "POINT_IN_TIME", "key": "submission_benchmark", "value": "llama31_8b", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1333}}
+:::MLLOG {"namespace": "", "time_ms": 1778312378485, "event_type": "POINT_IN_TIME", "key": "cache_clear", "value": true, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1336}}
+:::MLLOG {"namespace": "", "time_ms": 1778312378485, "event_type": "INTERVAL_START", "key": "init_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1337}}
+:::MLLOG {"namespace": "", "time_ms": 1778312726494, "event_type": "INTERVAL_END", "key": "init_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1629}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740045, "event_type": "INTERVAL_START", "key": "run_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1340}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740049, "event_type": "POINT_IN_TIME", "key": "seed", "value": 12228, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1341}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740049, "event_type": "POINT_IN_TIME", "key": "global_batch_size", "value": 32, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1343}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740049, "event_type": "POINT_IN_TIME", "key": "max_sequence_length", "value": 8192, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1344}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "max_steps", "value": 1200000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1345}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "gradient_accumulation_steps", "value": 2, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1346}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "eval_samples", "value": 1024, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1347}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "train_samples", "value": 38400000, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1348}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_name", "value": "adamw", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1350}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_base_learning_rate", "value": 0.001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1351}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_end_learning_rate", "value": 0.0001, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1352}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_1", "value": 0.9, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1353}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_adamw_beta_2", "value": 0.95, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1354}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_adamw_epsilon", "value": 1e-05, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1355}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_adamw_weight_decay", "value": 0.1, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1356}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1357}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740050, "event_type": "POINT_IN_TIME", "key": "num_warmup_steps", "value": 128, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1358}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740051, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_steps", "value": 1199872, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1359}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740051, "event_type": "POINT_IN_TIME", "key": "opt_learning_rate_decay_schedule", "value": "cosine with linear warmup", "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1360}}
+:::MLLOG {"namespace": "", "time_ms": 1778312740051, "event_type": "POINT_IN_TIME", "key": "opt_gradient_clip_norm", "value": 1.0, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1361}}
+:::MLLOG {"namespace": "", "time_ms": 1778313057094, "event_type": "INTERVAL_START", "key": "epoch_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1529, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778313057095, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1530, "samples_count": 0}}
+:::MLLOG {"namespace": "", "time_ms": 1778313872567, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778313872567, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778313917470, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 5.736083984375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778313917471, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778313917472, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 12288}}
+:::MLLOG {"namespace": "", "time_ms": 1778314572849, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778314572850, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778314602523, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.584716796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778314602524, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778314602525, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 24576}}
+:::MLLOG {"namespace": "", "time_ms": 1778315258897, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778315258898, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778315288494, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 4.114501953125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778315288495, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778315288496, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 36864}}
+:::MLLOG {"namespace": "", "time_ms": 1778315946776, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778315946777, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778315976384, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.906005859375, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778315976385, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778315976386, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 49152}}
+:::MLLOG {"namespace": "", "time_ms": 1778316632177, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778316632178, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778316661800, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.76513671875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778316661800, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778316661801, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 61440}}
+:::MLLOG {"namespace": "", "time_ms": 1778317318705, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778317318706, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778317348421, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.6568603515625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778317348421, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778317348422, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 73728}}
+:::MLLOG {"namespace": "", "time_ms": 1778318007246, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778318007246, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778318036837, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.5897216796875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778318036838, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778318036839, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 86016}}
+:::MLLOG {"namespace": "", "time_ms": 1778318691769, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778318691770, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778318721376, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.52587890625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778318721377, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778318721377, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 98304}}
+:::MLLOG {"namespace": "", "time_ms": 1778319374807, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778319374808, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778319404256, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.473388671875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778319404257, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778319404258, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 110592}}
+:::MLLOG {"namespace": "", "time_ms": 1778320058613, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778320058613, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778320087986, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.4307861328125, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778320087987, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778320087988, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 122880}}
+:::MLLOG {"namespace": "", "time_ms": 1778320742022, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778320742022, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778320771659, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3931884765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778320771660, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778320771660, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 135168}}
+:::MLLOG {"namespace": "", "time_ms": 1778321426019, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778321426019, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778321455724, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3629150390625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778321455725, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778321455726, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 147456}}
+:::MLLOG {"namespace": "", "time_ms": 1778322114634, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778322114634, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778322144126, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3377685546875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778322144127, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778322144127, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 159744}}
+:::MLLOG {"namespace": "", "time_ms": 1778322801727, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778322801728, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778322831371, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.3150634765625, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778322831372, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778322831372, "event_type": "INTERVAL_START", "key": "block_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1654, "samples_count": 172032}}
+:::MLLOG {"namespace": "", "time_ms": 1778323487126, "event_type": "INTERVAL_END", "key": "block_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1616, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778323487126, "event_type": "INTERVAL_START", "key": "eval_start", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1617, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778323516691, "event_type": "POINT_IN_TIME", "key": "eval_accuracy", "value": 3.2889404296875, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1637, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778323516691, "event_type": "INTERVAL_END", "key": "eval_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1638, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778323516692, "event_type": "INTERVAL_END", "key": "epoch_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1646, "samples_count": 184320}}
+:::MLLOG {"namespace": "", "time_ms": 1778323516692, "event_type": "INTERVAL_END", "key": "run_stop", "value": null, "metadata": {"file": "tinygrad3/examples/mlperf/model_train.py", "lineno": 1647, "status": "success"}}

examples/mlperf/training_submission_v6.0/tinycorp/systems/tinybox_8xMI300X.json

@@ -1,38 +0,0 @@
-{
-    "submitter": "tinycorp",
-    "division": "closed",
-    "status": "Available on-premise",
-    "system_name": "tinybox 8xMI300X",
-    "number_of_nodes": "1",
-    "host_processors_per_node": "2",
-    "host_processor_model_name": "AMD EPYC 9354",
-    "host_processor_core_count": "32",
-    "host_processor_vcpu_count": "64",
-    "host_processor_frequency": "",
-    "host_processor_caches": "",
-    "host_processor_interconnect": "",
-    "host_memory_capacity": "2304GB",
-    "host_storage_type": "NVMe SSD",
-    "host_storage_capacity": "3x 4TB raid array",
-    "host_networking": "",
-    "host_networking_topology": "",
-    "host_memory_configuration": "24x 96GB DDR5",
-    "accelerators_per_node": "8",
-    "accelerator_model_name": "AMD Instinct MI300X 192GB HBM3",
-    "accelerator_host_interconnect": "PCIe 5.0 x16",
-    "accelerator_frequency": "",
-    "accelerator_on-chip_memories": "",
-    "accelerator_memory_configuration": "HBM3",
-    "accelerator_memory_capacity": "192GB",
-    "accelerator_interconnect": "",
-    "accelerator_interconnect_topology": "",
-    "cooling": "air",
-    "hw_notes": "",
-    "framework": "tinygrad, branch mlperf_training_v5.0",
-    "other_software_stack": {
-        "python": "3.10.16",
-        "ROCm": "3.0.0+94441cb"
-    },
-    "operating_system": "Ubuntu 24.04.1 LTS",
-    "sw_notes": ""
-  }

examples/mlperf/training_submission_v6.0/tinycorp/systems/tinybox_8xMI350X.json

@@ -34,5 +34,5 @@
         "ROCm": "7.1.1"
     },
     "operating_system": "Ubuntu 24.04.3 LTS",
-    "sw_notes": ""
+    "sw_notes": "tinygrad @ 026688f03f84a75ec3fef034bcba916bf8f8bdc6"
   }

examples/mlperf/training_submission_v6.0/tinycorp/systems/tinybox_green.json

@@ -1,38 +0,0 @@
-{
-  "submitter": "tinycorp",
-  "division": "closed",
-  "status": "Available on-premise",
-  "system_name": "tinybox green",
-  "number_of_nodes": "1",
-  "host_processors_per_node": "1",
-  "host_processor_model_name": "AMD EPYC 7532",
-  "host_processor_core_count": "32",
-  "host_processor_vcpu_count": "64",
-  "host_processor_frequency": "",
-  "host_processor_caches": "",
-  "host_processor_interconnect": "",
-  "host_memory_capacity": "128GB",
-  "host_storage_type": "NVMe SSD",
-  "host_storage_capacity": "4 TB raid array + 1 TB boot",
-  "host_networking": "",
-  "host_networking_topology": "",
-  "host_memory_configuration": "8x 16GB DDR4",
-  "accelerators_per_node": "6",
-  "accelerator_model_name": "NVIDIA GeForce RTX 4090",
-  "accelerator_host_interconnect": "PCIe 4.0 x16",
-  "accelerator_frequency": "",
-  "accelerator_on-chip_memories": "",
-  "accelerator_memory_configuration": "GDDR6X",
-  "accelerator_memory_capacity": "24GB",
-  "accelerator_interconnect": "",
-  "accelerator_interconnect_topology": "",
-  "cooling": "air",
-  "hw_notes": "",
-  "framework": "tinygrad, branch mlperf_training_v5.0",
-  "other_software_stack": {
-    "python": "3.10.12",
-    "CUDA": "12.4"
-  },
-  "operating_system": "Ubuntu 22.04.4",
-  "sw_notes": ""
-}

examples/mlperf/training_submission_v6.0/tinycorp/systems/tinybox_red.json

@@ -1,37 +0,0 @@
-{
-  "submitter": "tinycorp",
-  "division": "closed",
-  "status": "Available on-premise",
-  "system_name": "tinybox red",
-  "number_of_nodes": "1",
-  "host_processors_per_node": "1",
-  "host_processor_model_name": "AMD EPYC 7532",
-  "host_processor_core_count": "32",
-  "host_processor_vcpu_count": "64",
-  "host_processor_frequency": "",
-  "host_processor_caches": "",
-  "host_processor_interconnect": "",
-  "host_memory_capacity": "128GB",
-  "host_storage_type": "NVMe SSD",
-  "host_storage_capacity": "4 TB raid array + 1 TB boot",
-  "host_networking": "",
-  "host_networking_topology": "",
-  "host_memory_configuration": "8x 16GB DDR4",
-  "accelerators_per_node": "6",
-  "accelerator_model_name": "AMD Radeon RX 7900 XTX",
-  "accelerator_host_interconnect": "PCIe 4.0 x16",
-  "accelerator_frequency": "",
-  "accelerator_on-chip_memories": "",
-  "accelerator_memory_configuration": "GDDR6",
-  "accelerator_memory_capacity": "24GB",
-  "accelerator_interconnect": "",
-  "accelerator_interconnect_topology": "",
-  "cooling": "air",
-  "hw_notes": "",
-  "framework": "tinygrad, branch mlperf_training_v5.0",
-  "other_software_stack": {
-    "python": "3.10.12"
-  },
-  "operating_system": "Ubuntu 22.04.4",
-  "sw_notes": ""
-}

examples/openpilot/compile3.py

@@ -4,7 +4,7 @@ if "JIT_BATCH_SIZE" not in os.environ: os.environ["JIT_BATCH_SIZE"] = "0"
 
 from tinygrad import fetch, Tensor, TinyJit, Context, GlobalCounters, Device, dtypes
 from tinygrad.helpers import DEBUG, getenv
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.uop.ops import Ops
 from tinygrad.nn.onnx import OnnxRunner
 
 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"
@@ -35,7 +35,11 @@ def compile(onnx_file):
       ret = run_onnx_jit(**inputs).numpy()
     # copy i == 1 so use of JITBEAM is okay
     if i == 1: test_val = np.copy(ret)
-  print(f"captured {len(run_onnx_jit.captured.jit_cache)} kernels")
+  # iterate kernel CALLs in the captured LINEAR UOp; toposort descends into batched graph CUSTOM_FUNCTIONs
+  kernel_asts = {Ops.PROGRAM}
+  kernel_calls = [u for u in run_onnx_jit.captured.linear.toposort(gate=lambda x: x.op not in kernel_asts)
+                  if u.op is Ops.CALL and u.src[0].op in kernel_asts]
+  print(f"captured {len(kernel_calls)} kernels")
   np.testing.assert_equal(test_val, ret, "JIT run failed")
   print("jit run validated")
 
@@ -43,13 +47,14 @@ def compile(onnx_file):
   kernel_count = 0
   read_image_count = 0
   gated_read_image_count = 0
-  for ei in run_onnx_jit.captured.jit_cache:
-    if isinstance(ei.prg, CompiledRunner):
-      kernel_count += 1
-      read_image_count += ei.prg.p.src.count("read_image")
-      gated_read_image_count += ei.prg.p.src.count("?read_image")
-      for v in [m.group(1) for m in re.finditer(r'(val\d+)\s*=\s*read_imagef\(', ei.prg.p.src)]:
-        if len(re.findall(fr'[\?\:]{v}\.[xyzw]', ei.prg.p.src)) > 0: gated_read_image_count += 1
+  for call in kernel_calls:
+    _, _, _, source, _ = call.src[0].src
+    src = source.arg
+    kernel_count += 1
+    read_image_count += src.count("read_image")
+    gated_read_image_count += src.count("?read_image")
+    for v in [m.group(1) for m in re.finditer(r'(val\d+)\s*=\s*read_imagef\(', src)]:
+      if len(re.findall(fr'[\?\:]{v}\.[xyzw]', src)) > 0: gated_read_image_count += 1
   print(f"{kernel_count=},  {read_image_count=}, {gated_read_image_count=}")
   if (allowed_kernel_count:=getenv("ALLOWED_KERNEL_COUNT", -1)) != -1:
     assert kernel_count == allowed_kernel_count, f"different kernels! {kernel_count=}, {allowed_kernel_count=}"
@@ -128,14 +133,20 @@ def bench(run, inputs):
       run(**inputs).numpy()
 
 if __name__ == "__main__":
-  onnx_file = fetch(OPENPILOT_MODEL)
-  inputs, outputs = compile(onnx_file)
+  if getenv("RUN_PICKLE"):
+    with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)
+    inputs = {name: Tensor(Tensor.randn(*[int(s) for s in view.src[1].arg], dtype=dtype).numpy(), device=device)
+              for name, (view, _vars, dtype, device) in zip(pickle_loaded.captured.expected_names, pickle_loaded.captured.expected_input_info)}
+    test_vs_compile(pickle_loaded, inputs)
+  else:
+    onnx_file = fetch(OPENPILOT_MODEL)
+    inputs, outputs = compile(onnx_file)
 
-  with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)
+    with open(OUTPUT, "rb") as f: pickle_loaded = pickle.load(f)
 
-  test_vs_compile(pickle_loaded, inputs, outputs)
-  if getenv("SELFTEST"):
-    test_vs_onnx(inputs, outputs, onnx_file, 1e-4)
+    test_vs_compile(pickle_loaded, inputs, outputs)
+    if getenv("SELFTEST"):
+      test_vs_onnx(inputs, outputs, onnx_file, 1e-4)
 
   if getenv("BENCHMARK_LOG", ""):
     bench(pickle_loaded, inputs)

examples/tools/gpuburn.py

@@ -1,7 +1,6 @@
 from tinygrad import Tensor, Device, TinyJit, dtypes
-from tinygrad.helpers import getenv
 
-GPUS = getenv("GPUS", 4) # TODO: expose a way in tinygrad to access this
+GPUS = Device[Device.DEFAULT].count()
 N = 6144
 
 @TinyJit

examples/webgpu/stable_diffusion/compile.py

@@ -111,19 +111,19 @@ if __name__ == "__main__":
     return code
 
   def compile_step(model, step: Step):
-    run, special_names = jit_model(step, *step.input)
-    functions, statements, bufs, _ = compile_net(run, special_names)
+    linear, output_bufs = jit_model(step, *step.input)
+    functions, statements, bufs, _ = compile_net(linear, output_bufs)
     state = get_state_dict(model)
-    weights = {id(x.uop.base.realized): name for name, x in state.items()}
+    weights = {(id(b), b.offset, b.size, b.dtype): name for name, x in state.items() if (b:=x.uop.base.realized) is not None}
     kernel_code = '\n\n'.join([f"const {key} = `{fixup_code(code, key)}`;" for key, code in functions.items()])
     kernel_names = ', '.join([name for (name, _, _, _) in statements])
-    input_names = [name for _,name in special_names.items() if "input" in name]
-    output_names = [name for _,name in special_names.items() if "output" in name]
+    input_names = [f"input{i}" for i in range(len(step.input))]
+    output_names = [f"output{i}" for i in range(len(output_bufs))]
     input_buf_types = [dtype_to_js_type(bufs[inp_name][1]) for inp_name in input_names]
     output_buf_types = [dtype_to_js_type(bufs[out_name][1]) for out_name in output_names]
     kernel_calls = '\n        '.join([f"addComputePass(device, commandEncoder, piplines[{i}], [{', '.join(args)}], {global_size});" for i, (_name, args, global_size, _local_size) in enumerate(statements) ])
     exported_bufs =  '\n    '.join([f"const {name} = " + (f"createEmptyBuf(device, {size});" if _key not in weights else f"createWeightBuf(device, {size}, getTensorBuffer(safetensor, metadata['{weights[_key]}'], '{weights[_key]}'))") + ";"  for name,(size,dtype,_key) in bufs.items()])
-    gpu_write_bufs =  '\n    '.join([f"const gpuWriteBuffer{i} = device.createBuffer({{size:input{i}.size, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE }});" for i,(_,value) in enumerate(special_names.items()) if "output" not in value])
+    gpu_write_bufs =  '\n    '.join([f"const gpuWriteBuffer{i} = device.createBuffer({{size:input{i}.size, usage: GPUBufferUsage.COPY_SRC | GPUBufferUsage.MAP_WRITE }});" for i in range(len(input_names))])
     input_writer = '\n    '.join([f"await gpuWriteBuffer{i}.mapAsync(GPUMapMode.WRITE);\n    new {input_buf_types[i]}(gpuWriteBuffer{i}.getMappedRange()).set(" + f'data{i});' + f"\n    gpuWriteBuffer{i}.unmap();\ncommandEncoder.copyBufferToBuffer(gpuWriteBuffer{i}, 0, input{i}, 0, gpuWriteBuffer{i}.size);"  for i,_ in enumerate(input_names)])
     return f"""\n    var {step.name} = function() {{
 
@@ -141,7 +141,7 @@ if __name__ == "__main__":
         const kernels = [{kernel_names}];
         const piplines = await Promise.all(kernels.map(name => device.createComputePipelineAsync({{layout: "auto", compute: {{ module: device.createShaderModule({{ code: name }}), entryPoint: "main" }}}})));
 
-        return async ({",".join([f'data{i}' for i,(k,v) in enumerate(special_names.items()) if v != "output0"])}) => {{
+        return async ({",".join([f'data{i}' for i in range(len(input_names))])}) => {{
             const commandEncoder = device.createCommandEncoder();
 
             {input_writer}

extra/amdpci/am_smi.py

@@ -64,7 +64,7 @@ def get_bar0_size(pcibus):
 
 class AMSMI(AMDev):
   def __init__(self, pcibus, vram_bar:MMIOInterface, doorbell_bar:MMIOInterface, mmio_bar:MMIOInterface):
-    self.pcibus = pcibus
+    self.pcibus, self.devfmt = pcibus, pcibus
     self.vram, self.doorbell64, self.mmio = vram_bar, doorbell_bar, mmio_bar
     self.pci_state = self.read_pci_state()
     if self.pci_state == "D0": self._init_from_d0()
@@ -91,6 +91,7 @@ class SMICtx:
     self.prev_lines_cnt = 0
     self.prev_terminal_width = 0
     self.prev_terminal_height = 0
+    self.prev_metrics = {}
 
     remove_parts = ["Advanced Micro Devices, Inc. [AMD/ATI]", "VGA compatible controller:", "Processing accelerators:"]
     lspci = subprocess.check_output(["lspci"]).decode("utf-8").splitlines()
@@ -235,6 +236,29 @@ class SMICtx:
       case (13,0,12): return self._smuq10_round(metrics.SocketPower), self._smuq10_round(metrics.SocketPowerLimit)
       case _: return metrics.SmuMetrics.AverageSocketPower, metrics.SmuMetrics.dGPU_W_MAX
 
+  def get_throttle_info(self, dev, metrics):
+    match dev.ip_ver[am.MP1_HWIP]:
+      case (13,0,6)|(13,0,12):
+        throttle_fields = [('ProchotResidencyAcc', 'Prochot'), ('PptResidencyAcc', 'PPT'),
+                           ('SocketThmResidencyAcc', 'Socket Thm'), ('VrThmResidencyAcc', 'VR Thm'), ('HbmThmResidencyAcc', 'HBM Thm')]
+        prev = self.prev_metrics.get(dev.pcibus)
+        active = []
+        if prev is not None:
+          acc_delta = metrics.AccumulationCounter - prev.AccumulationCounter
+          if acc_delta > 0:
+            for field, name in throttle_fields:
+              delta = getattr(metrics, field) - getattr(prev, field)
+              if delta > 0 and (pct := min(100, (delta * 100 + acc_delta // 2) // acc_delta)) > 0: active.append((name, pct))
+        return active
+      case _:
+        smu_mod = dev.smu.smu_mod
+        throttler_names = {getattr(smu_mod, a): a[len('THROTTLER_'):-len('_BIT')]
+                           for a in dir(smu_mod) if a.startswith('THROTTLER_') and a.endswith('_BIT')}
+        active = []
+        for i, pct in enumerate(metrics.SmuMetrics.ThrottlingPercentage):
+          if pct > 0: active.append((throttler_names.get(i, f"UNK_{i}"), int(pct)))
+        return active
+
   def get_mem_usage(self, dev):
     usage = 0
     pt_stack = [dev.mm.root_page_table]
@@ -281,6 +305,13 @@ class SMICtx:
                     + [f"MEM Activity {draw_bar(self.get_mem_activity(dev, metrics) / 100, activity_line_width)}"] \
                     + [f"MEM Usage    {draw_bar(mem_used / mem_total, activity_line_width, opt_text=mem_fmt)}"] \
 
+      throttle_info = self.get_throttle_info(dev, metrics)
+      if throttle_info:
+        throttle_text = colored(', '.join(f"{name} {pct}%" for name, pct in throttle_info), "red")
+      else:
+        throttle_text = colored("None", "green")
+      activity_line += [f"Throttle     {throttle_text}" + " " * (activity_line_width + 2)]
+
       temps_data, temps_data_compact = self.get_temps(dev, metrics), self.get_temps(dev, metrics, compact=True)
       temps_table = ["=== Temps (°C) ==="] + [f"{name:<16}: {color_temp(val)}" for name, val in temps_data.items()]
       temps_table_compact = ["Temps (°C):" + '/'.join([f"{color_temp(val)} {name}" for name, val in temps_data_compact.items()])]
@@ -324,6 +355,8 @@ class SMICtx:
 
       dev_content.append(device_line + activity_line + same_line([temps_table, power_table, frequency_table]))
 
+    self.prev_metrics = {dev.pcibus: m for dev, m in dev_metrics.items() if m is not None}
+
     raw_text = 'AM Monitor'.center(terminal_width) + "\n" + "=" * terminal_width + "\n\n"
     for i in range(0, len(dev_content), 2):
       if i + 1 < len(dev_content): raw_text += '\n'.join(same_line([dev_content[i], dev_content[i+1]], split=padding))

extra/amdpci/headers/amdgpu_irq.h

@@ -28,15 +28,7 @@
 // #include "soc15_ih_clientid.h"
 // #include "amdgpu_ih.h"
 
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 #define AMDGPU_MAX_IRQ_SRC_ID		0x100
 #define AMDGPU_MAX_IRQ_CLIENT_ID	0x100

extra/amdpci/headers/amdgpu_smu.h

@@ -22,15 +22,7 @@
 #ifndef __AMDGPU_SMU_H__
 #define __AMDGPU_SMU_H__
 
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 #define SMU_THERMAL_MINIMUM_ALERT_TEMP		0
 #define SMU_THERMAL_MAXIMUM_ALERT_TEMP		255

extra/amdpci/headers/amdgpu_ucode.h

@@ -24,15 +24,7 @@
 #define __AMDGPU_UCODE_H__
 
 // #include "amdgpu_socbb.h"
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 struct common_firmware_header {
 	uint32_t size_bytes; /* size of the entire header+image(s) in bytes */

extra/export_model.py

@@ -1,47 +1,50 @@
 from typing import Tuple, Dict, List, Optional
 from tinygrad.dtype import DType, dtypes
-from tinygrad.renderer import ProgramSpec
 from tinygrad.tensor import Tensor
-from tinygrad.device import Device
+from tinygrad.device import Device, Buffer
 from tinygrad.engine.jit import TinyJit
 from tinygrad.nn.state import get_state_dict
-from tinygrad.helpers import Context, to_mv
-from tinygrad.uop.ops import Ops
+from tinygrad.helpers import Context, to_mv, prod
+from tinygrad.uop.ops import Ops, UOp
+from tinygrad.codegen import to_program
 import json
 from collections import OrderedDict
 
 EXPORT_SUPPORTED_DEVICE = ["WEBGPU", "CPU", "CUDA", "CL"]
 
-def compile_net(run:TinyJit, special_names:Dict[int,str]) -> Tuple[Dict[str,str],List[Tuple[str,List[str],List[int]]],Dict[str,Tuple[int,DType,int]],Dict[str,Tensor]]:
-  # memory-planned subbuffers can have multiple Buffer objects for the same memory region
-  canon, _seen = {}, {}
-  for ji in run.jit_cache:
-    for b in ji.bufs:
-      if b is not None: canon[id(b)] = _seen.setdefault((id(b.base._buf), b.offset, b.size, b.dtype), b)
-  special_names = {id(canon[k]): v for k, v in special_names.items() if k in canon}
+_KERNEL_ASTS = {Ops.SINK, Ops.PROGRAM}
+def iter_kernel_calls(linear:UOp):
+  """Yield kernel CALLs from a LINEAR UOp. Toposort descends naturally into CUSTOM_FUNCTION graph batches; gate stops at kernel ASTs."""
+  return (u for u in linear.toposort(gate=lambda x: x.op not in _KERNEL_ASTS) if u.op is Ops.CALL and u.src[0].op in _KERNEL_ASTS)
 
-  functions, bufs, bufs_to_save, statements, bufnum = {}, {}, {}, [], 0
-  for ji in run.jit_cache:
-    fxn: ProgramSpec = ji.prg.p
-    functions[fxn.function_name] = fxn.src   # NOTE: this assumes all with the same name are the same
-    cargs = []
-    for i,arg in enumerate(ji.bufs):
-      arg = canon[id(arg)]
-      key = id(arg)
-      if key not in bufs:
-        if key in special_names:
-          bufs[key] = (special_names[key], arg.size*arg.dtype.itemsize, arg.dtype, key)
-        else:
-          bufs[key] = (f"buf_{bufnum}", arg.size*arg.dtype.itemsize, arg.dtype, key)
-          bufnum += 1
-          if i > 0: bufs_to_save[bufs[key][0]] = arg   # if first usage of a buffer is not an output, and it's not a special name
-      cargs.append(bufs[key][0])
-    cargs += [var for var in fxn.vars if getattr(var, "op", None) is Ops.DEFINE_VAR] # symbolic vars; is it necessary or sufficient to check for DEFINE_VAR?
-    statements.append((fxn.function_name, cargs, fxn.global_size, fxn.local_size))
+def compile_net(linear:UOp, output_bufs:List[Buffer]) -> Tuple[Dict[str,str], List, Dict[str,Tuple[int,DType,int]], Dict[str,Buffer]]:
+  output_name = {id(b): f"output{i}" for i, b in enumerate(output_bufs)}
+  functions, bufs, bufs_to_save, statements, n = {}, {}, {}, [], 0
 
-  return functions, statements, {name:(size, dtype, key) for (name,size,dtype,key) in bufs.values()}, bufs_to_save
+  def name_of(bu:UOp, is_out:bool) -> str:
+    nonlocal n
+    if bu.op is Ops.PARAM: key, name, size = ("in", bu.arg), f"input{bu.arg}", prod(bu.shape)*bu.dtype.itemsize
+    else:
+      b = bu.buffer
+      key, size = (id(b.base), b.offset, b.size, b.dtype), b.size*b.dtype.itemsize
+      if key in bufs: return bufs[key][0]
+      if (name:=output_name.get(id(b))) is None:
+        name, n = f"buf_{n}", n+1
+        if not is_out: bufs_to_save[name] = b
+    bufs[key] = (name, size, bu.dtype, key)
+    return name
 
-def jit_model(model, *args) -> Tuple[TinyJit,Dict[int,str]]:
+  for call in iter_kernel_calls(linear):
+    arg_uops = [b for b in call.src[1:] if b.op is not Ops.BIND]
+    prg = to_program(call.src[0], Device[arg_uops[0].device].renderer)
+    info = prg.arg
+    functions[info.function_name] = prg.src[3].arg
+    cargs = [name_of(bu, i == 0) for i, bu in enumerate(arg_uops)] + [v for v in info.vars if v.op is Ops.DEFINE_VAR]
+    statements.append((info.function_name, cargs, info.global_size, info.local_size))
+
+  return functions, statements, {name:(size, dtype, key) for name, size, dtype, key in bufs.values()}, bufs_to_save
+
+def jit_model(model, *args) -> Tuple[UOp, List[Buffer]]:
   assert hasattr(model, "forward") or callable(model), "model needs a forward function"
   @TinyJit
   def run(*x):
@@ -50,20 +53,10 @@ def jit_model(model, *args) -> Tuple[TinyJit,Dict[int,str]]:
     out = [out] if isinstance(out, Tensor) else out
     return [o.realize() for o in out]
 
-  # twice to run the JIT
+  # run twice to trigger JIT capture
   for _ in range(2): the_output = run(*args)
-  special_names = {}
-
-  # hack to put the inputs back
-  for (j,i),idx in run.input_replace.items():
-    realized_input = args[idx].uop.base.realized
-    run.jit_cache[j].bufs[i] = realized_input
-    special_names[id(realized_input)] = f'input{idx}'
-
-  # TODO: fetch this from the jit in self.input_replace and self.ret (hint: use get_parameters on self.ret)
-  for i, output in enumerate(the_output):
-    special_names[id(output.uop.base.realized)] = f'output{i}'
-  return run, special_names
+  assert run.captured is not None
+  return run.captured.linear, [o.uop.base.realized for o in the_output]
 
 def export_model_clang(functions:Dict[str,str], statements:Dict[str,Tuple[str,int,int]], bufs:Dict[str,Tuple[str,int,int]],
   bufs_to_save:Dict[str,Tensor], input_names:List[str], output_names:List[str], weight_names={}, model_name="model", symbolic_vars={}, wasm=False) -> str:
@@ -249,12 +242,12 @@ def export_model(model, target:str, *inputs, model_name: Optional[str] = "model"
   assert Device.DEFAULT in EXPORT_SUPPORTED_DEVICE, f"only {', '.join(EXPORT_SUPPORTED_DEVICE)} are supported"
 
   # NOTE: CPU_COUNT=1, since export does not support threading
-  with Context(JIT=2, CPU_COUNT=1): run,special_names = jit_model(model, *inputs)
-  functions, statements, bufs, bufs_to_save = compile_net(run, special_names)
+  with Context(JIT=2, CPU_COUNT=1): linear, output_bufs = jit_model(model, *inputs)
+  functions, statements, bufs, bufs_to_save = compile_net(linear, output_bufs)
   state = get_state_dict(model)
-  weight_names = {id(x.uop.base.realized): name for name, x in state.items()}
-  input_names = [name for _,name in special_names.items() if "input" in name]
-  output_names = [name for _,name in special_names.items() if "output" in name]
+  weight_names = {(id(b), b.offset, b.size, b.dtype): name for name, x in state.items() if (b:=x.uop.base.realized) is not None}
+  input_names = [f"input{i}" for i in range(len(inputs))]
+  output_names = [f"output{i}" for i in range(len(output_bufs))]
 
   # handle symbolic variables; TODO: refactor to fix some of this stuff upstream in tinygrad
   symbolic_vars = OrderedDict()

extra/gemm/amd_asm_matmul.py

@@ -13,7 +13,7 @@ from tinygrad import Tensor, Device, Context, GlobalCounters
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.helpers import getenv, colored
 from tinygrad.dtype import dtypes, AddrSpace
-from tinygrad.engine.realize import Estimates
+from tinygrad.engine.realize import Estimates, run_linear
 from tinygrad.renderer.amd.dsl import s, v, VCC_LO, NULL
 from tinygrad.runtime.autogen.amd.rdna3.ins import *
 
@@ -167,7 +167,7 @@ PREFETCH_LOADS = [(V_LDS_A_DATA[4+2*i], V_LDS_A_DATA[4+2*i+1], V_GLOBAL_B_ADDR,
 # =============================================================================
 
 class Kernel:
-  def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
+  def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
   def label(self, name): self.labels[name] = self.pos
 
   def emit(self, inst, target=None):
@@ -196,10 +196,10 @@ class Kernel:
 # Kernel builder
 # =============================================================================
 
-def build_kernel(N, arch='gfx1100'):
+def build_kernel(N):
   assert N % 128 == 0, f"N must be a multiple of 128 (tile size), got {N}"
   assert N >= 256, f"N must be >= 256 (prefetch pipeline requires at least 2 K-blocks), got {N}"
-  k = Kernel(arch)
+  k = Kernel()
 
   # ===========================================================================
   # PROLOGUE: Load kernel arguments, compute tile coordinates and addresses
@@ -443,7 +443,7 @@ def test_matmul():
   dev = Device[Device.DEFAULT]
   print(f"Device arch: {dev.renderer.target.arch}")
 
-  insts = build_kernel(N, dev.renderer.target.arch)
+  insts = build_kernel(N)
 
   rng = np.random.default_rng(42)
   a = Tensor(rng.random((N, N), dtype=np.float32) - 0.5)
@@ -463,11 +463,14 @@ def test_matmul():
                                                                                   estimates=Estimates(ops=N*N*N*2, mem=N*N*4*3)))
     return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
   c = Tensor.custom_kernel(a, b, c, fxn=asm_kernel)[2]
-  ei = c.schedule()[0].lower()
+  linear = c.schedule_linear()
 
   ets = []
   with Context(DEBUG=2):
-    for _ in range(getenv("CNT", 5)): ets.append(ei.run(wait=True))
+    for _ in range(getenv("CNT", 5)):
+      start = GlobalCounters.time_sum_s
+      run_linear(linear)
+      ets.append(GlobalCounters.time_sum_s - start)
   print(f"REAL TFLOPS {N * N * N * 2 / min(ets) * 1e-12:.2f}")
 
   if getenv("VERIFY", 1):

extra/gemm/amd_matmul.py

@@ -1,31 +1,39 @@
 # kernel8_batched_gmem.s from https://seb-v.github.io/optimization/update/2025/01/20/Fast-GPU-Matrix-multiplication.html
 # sudo PATH=/opt/homebrew/Cellar/llvm/20.1.6/bin:$PATH AMD_LLVM=0 AMD=1 DEBUG=2 python3 extra/gemm/amd_matmul.py
 import pathlib
-from dataclasses import replace
 from tinygrad import Tensor, Device, Context, GlobalCounters
 from tinygrad.helpers import getenv
-from tinygrad.engine.realize import CompiledRunner, ExecItem, get_program
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from tinygrad.engine.realize import run_linear
 
 N = 4096
 run_count = 5
 
-if __name__ == "__main__":
-  ast = (Tensor.empty(N, N)@Tensor.empty(N, N)).schedule()[-1].ast
-  prg = get_program(ast, Device.default.renderer)
+def make_matmul_kernel(name:str, src:str, local_size:int):
+  def fxn(a:UOp, b:UOp, c:UOp) -> UOp:
+    threads = UOp.special(local_size, "lidx0")
+    wg_x = UOp.special(N//128, "gidx0")
+    wg_y = UOp.special(N//128, "gidx1")
+    sink = UOp.sink(a.base, b.base, c.base, threads, wg_x, wg_y, arg=KernelInfo(name, estimates=Estimates(ops=2*N**3, mem=3*N*N*4)))
+    lib = Device[Device.DEFAULT].compiler.compile_cached(src)
+    return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                                 UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+  return fxn
 
+if __name__ == "__main__":
   if getenv("ASM") == 1:
     src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel8_batched_gmem.s").read_text()
-    prgfast = replace(prg, name="kernel", src=src, global_size=[N//128, N//128, 1], local_size=[128, 1, 1])
+    name, local_size = "kernel", 128
   elif getenv("ASM") == -1:
     src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel3_registers.cpp").read_text()
-    prgfast = replace(prg, name="kernel3_registers", src=src, global_size=[N//128, N//128, 1], local_size=[256, 1, 1])
+    name, local_size = "kernel3_registers", 256
   elif getenv("ASM") == -2:
     src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel4_gmem_df.cpp").read_text()
-    prgfast = replace(prg, name="kernel4_gmem_db", src=src, global_size=[N//128, N//128, 1], local_size=[256, 1, 1])
+    name, local_size = "kernel4_gmem_db", 256
   else:
     src = (pathlib.Path(__file__).parent / "amd_seb" / "kernel5_lds_optim.cpp").read_text()
-    prgfast = replace(prg, name="kernel5_lds_optim", src=src, global_size=[N//128, N//128, 1], local_size=[128, 1, 1])
-  runner = CompiledRunner(prgfast)
+    name, local_size = "kernel5_lds_optim", 128
 
   a = Tensor.randn(N, N).realize()
   b = Tensor.randn(N, N).realize()
@@ -35,8 +43,8 @@ if __name__ == "__main__":
   with Context(DEBUG=2):
     for _ in range(run_count): tc = (a@b).realize()
 
+  linear = Tensor.custom_kernel(a, b, c, fxn=make_matmul_kernel(name, src, local_size))[2].schedule_linear()
   GlobalCounters.reset()
-  ei = ExecItem(ast, [a.uop.buffer, b.uop.buffer, c.uop.buffer], prg=runner)
   with Context(DEBUG=2):
-    for _ in range(run_count): ei.run(wait=True)
+    for _ in range(run_count): run_linear(linear)
   print(f"custom  {(c-tc).square().mean().item()}")

extra/gemm/cdna_asm_gemm.py

@@ -6,7 +6,7 @@ from tinygrad.renderer import Estimates
 from tinygrad.helpers import getenv, all_same, DEBUG
 from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
 from tinygrad.runtime.autogen.amd.cdna.ins import *
-from examples.mlperf.models.flat_llama import FP8_DTYPE, FP8_GRAD_DTYPE, matmul, quantize_fp8
+from examples.mlperf.models.flat_llama import FP8_DTYPE, FP8_GRAD_DTYPE, quantize_fp8
 
 # ** CDNA4 assembly gemm
 
@@ -2628,20 +2628,24 @@ def custom_asm_gemm(C:UOp, A:UOp, B:UOp, dname:str) -> UOp:
 # ** FP8 GEMM custom kernel
 
 @functools.cache
-def custom_hk_fp8_gemm(C:UOp, A:UOp, B:UOp, S:UOp, dname:str) -> UOp:
-  # A is (batch, M, K), B is (N, K) transposed, S is combined scale (scalar float)
+def custom_hk_fp8_gemm(C:UOp, A:UOp, B:UOp, *args:UOp, dname:str, scale_mode:int=3) -> UOp:
+  # scale_mode: 0=no scale, 1=x only, 2=w only, 3=both
+  n_scales = (1 if scale_mode & 1 else 0) + (1 if scale_mode & 2 else 0)
+  scales, extra = args[:n_scales], args[n_scales:]
   M, K = A.shape[0]*A.shape[1], A.shape[2]
   N, K2 = B.shape[(1 if B.ndim == 3 else 0):]
   assert K == K2, f"{A.shape} {B.shape}"
   block_size = 256
   threads = UOp.special(64 * 8, "lidx0")
   workgroups = UOp.special((M // block_size) * (N // block_size), "gidx0")
-  sink = UOp.sink(C.base, A.base, B.base, S.base, threads, workgroups,
+  sink_inputs = (C.base, A.base, B.base) + tuple(s.base for s in scales) + (threads, workgroups)
+  sink = UOp.sink(*sink_inputs,
                   arg=KernelInfo(f"hk_fp8_gemm_{M}_{N}_{K}", estimates=Estimates(ops=2*M*N*K, mem=(M*K+N*K)*A.dtype.itemsize+M*N*C.dtype.itemsize)))
   kittens_path = pathlib.Path(__file__).parent.parent/"thunder"/"amd"
   src = (kittens_path/"gemm_fp8.cpp").read_text()
   lib = HIPCCCompiler("gfx950", [f"-I{(kittens_path/'include').as_posix()}", "-std=c++20", "-DKITTENS_CDNA4", "-ffast-math",
-                                 "-DHIP_ENABLE_WARP_SYNC_BUILTINS", f"-DGEMM_M={M}", f"-DGEMM_N={N}", f"-DGEMM_K={K}"]).compile_cached(src)
+                                 "-DHIP_ENABLE_WARP_SYNC_BUILTINS", f"-DGEMM_M={M}", f"-DGEMM_N={N}", f"-DGEMM_K={K}",
+                                 f"-DSCALE_MODE={scale_mode}"]).compile_cached(src)
   return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)), UOp(Ops.SOURCE, arg=src),
                                UOp(Ops.BINARY, arg=lib)))
 
@@ -2698,19 +2702,37 @@ def custom_uop_gemm(C:UOp, A:UOp, B:UOp) -> UOp:
 
 def custom_gemm_bw(gradient:UOp, kernel:UOp):
   inputs = kernel.src[1:]
-  # fp8 scaled gemm has 4 inputs (out, a, b, scale), others have 3 (out, a, b)
-  if len(inputs) == 4:
-    out, a, b, scale = inputs
-    a_t, b_t, g_t, s_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device), Tensor(scale, device=a.device)
+  if inputs[1].dtype == FP8_DTYPE:
+    grad_amax_state = inputs[5] if len(inputs) == 6 else None
+    out, a, b, s_x, s_w = inputs[:5]
+    a_t, b_t, g_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device)
+    s_x_t, s_w_t = Tensor(s_x, device=a.device), Tensor(s_w, device=a.device)
     g_t = g_t[:a.shape[0]]
-    # backward GEMMs in fp8 with scale applied inside kernel to prevent bf16 overflow
-    g_fp8, g_scale, _ = quantize_fp8(g_t)
-    bw_scale = g_scale * s_t
-    # dgrad: g_fp8 @ weight (asm_gemm computes a@b)
-    grad_a = asm_gemm(g_fp8, b_t, combined_scale=bw_scale)
-    # wgrad: g_fp8.T @ activation = (N, batch*seq) @ (batch*seq, K) → use permute to preserve sharding
-    grad_b = asm_gemm(g_fp8.permute(2, 0, 1).reshape(g_t.shape[-1], -1), a_t.reshape(-1, a_t.shape[-1]), combined_scale=bw_scale)
-    return (None, grad_a.uop, grad_b.uop, None)
+    from extra.llama_kernels.cast_amax import _grad_fp8_mailbox
+    from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+    gbase = gradient.base if hasattr(gradient, "base") else gradient
+    mailbox_entry = _grad_fp8_mailbox.pop(gbase, None) or _grad_fp8_mailbox.pop(gradient, None)
+    if mailbox_entry is not None:
+      g_fp8_u, inv_scale_u, _new_amax_u, store_effect = mailbox_entry
+      g_fp8 = Tensor(g_fp8_u, device=a.device)[:a.shape[0]]
+      g_scale = Tensor(inv_scale_u, device=a.device)
+    else:
+      assert grad_amax_state is not None, "fp8 matmul bwd needs either a mailbox entry or a grad_amax_state"
+      g_fp8, g_scale, _, store_effect = quantize_fp8_delayed(g_t, Tensor(grad_amax_state, device=a.device))
+    # dgrad: uses g_scale * x_scale * w_scale
+    grad_a = asm_gemm(g_fp8, b_t, x_scale=g_scale * s_x_t, w_scale=s_w_t)
+    # wgrad: no w_scale
+    g_fp8_2d = g_fp8.reshape(-1, g_fp8.shape[-1])
+    if getenv("FAST_FP8_TRANSPOSE", 0) and g_fp8_2d.shape[0] % 64 == 0 and g_fp8_2d.shape[1] % 64 == 0:
+      from extra.llama_kernels.fp8_transpose import fast_fp8_transpose
+      g_fp8_T = fast_fp8_transpose(g_fp8_2d)
+    else:
+      g_fp8_T = g_fp8.permute(2, 0, 1).reshape(g_t.shape[-1], -1)
+    grad_b = asm_gemm(g_fp8_T, a_t.reshape(-1, a_t.shape[-1]), x_scale=g_scale * s_x_t)
+    # Attach the delayed-amax store effect (if any) to grad_a so realizing grads commits the amax update.
+    ret = (None, grad_a.uop.after(store_effect), grad_b.uop, None, None)
+    if len(inputs) == 6: ret = ret + (None,)
+    return ret
   else:
     out, a, b = inputs
     assert all_same([gradient.device, a.device, b.device, out.device])
@@ -2725,7 +2747,7 @@ def custom_gemm_bw(gradient:UOp, kernel:UOp):
 
 # ** main gemm function
 
-def asm_gemm(a:Tensor, b:Tensor, combined_scale:Tensor|None=None) -> Tensor:
+def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=None, grad_amax_state:Tensor|None=None) -> Tensor:
   assert can_use_asm_gemm(a, b), f"{counters['todos'][-1]}"
   counters["used"] += 1
   unfold_batch = a.ndim == 3 and isinstance(a.device, tuple) and a.uop.axis == 2 and b.uop.axis == 0
@@ -2745,22 +2767,25 @@ def asm_gemm(a:Tensor, b:Tensor, combined_scale:Tensor|None=None) -> Tensor:
 
   if is_multi:
     if n_sharded:
-      out = Tensor(Tensor.invalid(batch, M, N//len(a.device), dtype=out_dtype, device=a.device).uop.multi(2), device=a.device)
+      out = Tensor(Tensor.invalids(batch, M, N//len(a.device), dtype=out_dtype, device=a.device).uop.multi(2), device=a.device)
     elif m_sharded:
-      out = Tensor(Tensor.invalid(batch, M, N, dtype=out_dtype, device=a.device).uop.multi(1), device=a.device)
+      out = Tensor(Tensor.invalids(batch, M, N, dtype=out_dtype, device=a.device).uop.multi(1), device=a.device)
     else:
-      out = Tensor(Tensor.invalid(batch//len(a.device) if a.uop.axis==0 else batch, M, N, dtype=out_dtype, device=a.device).uop.multi(0),
+      out = Tensor(Tensor.invalids(batch//len(a.device) if a.uop.axis==0 else batch, M, N, dtype=out_dtype, device=a.device).uop.multi(0),
                    device=a.device)
   else:
-    out = Tensor.invalid(batch, M, N, dtype=out_dtype, device=a.device)
+    out = Tensor.invalids(batch, M, N, dtype=out_dtype, device=a.device)
 
   renderer = Device[dname:=(a.device[0] if is_multi else a.device)].renderer
   dname, arch = dname.split(":")[0], renderer.target.arch
   if arch.startswith("gfx950") and getenv("USE_ASM", 1):
-    # fp8 gemm computes [email protected], with optional combined scale applied inside kernel before bf16 store
+    # fp8 gemm computes [email protected], kernel multiplies output by x_scale * w_scale before bf16 store
     if a.dtype == FP8_DTYPE:
-      scale = combined_scale if combined_scale is not None else Tensor(1.0, dtype=dtypes.float, device=a.device)
-      out = Tensor.custom_kernel(out, a, b.T, scale, fxn=functools.partial(custom_hk_fp8_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
+      scales = tuple(s for s in (x_scale, w_scale) if s is not None)
+      scale_mode = (1 if x_scale is not None else 0) | (2 if w_scale is not None else 0)
+      extra = [grad_amax_state] if grad_amax_state is not None else []
+      fxn = functools.partial(custom_hk_fp8_gemm, dname=dname, scale_mode=scale_mode)
+      out = Tensor.custom_kernel(out, a, b.T, *scales, *extra, fxn=fxn, grad_fxn=custom_gemm_bw)[0]
     else:
       out = Tensor.custom_kernel(out, a, b, fxn=functools.partial(custom_asm_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
   else:

extra/gemm/max_matmul.py

@@ -1,7 +1,6 @@
 import numpy as np, os
 from tinygrad.helpers import getenv, flat_mv
 from tinygrad import dtypes
-from tinygrad.engine.realize import get_program
 
 # for copied uops
 from tinygrad import dtypes

extra/gemm/rdna4_asm_matmul.py

@@ -4,7 +4,7 @@ from tinygrad import Tensor, Device, Context, GlobalCounters
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.helpers import getenv, colored
 from tinygrad.dtype import dtypes, AddrSpace
-from tinygrad.engine.realize import Estimates
+from tinygrad.engine.realize import Estimates, run_linear
 from tinygrad.renderer.amd.dsl import s, v, VCC_LO, NULL, src, ttmp
 from tinygrad.runtime.autogen.amd.rdna4.ins import *
 
@@ -225,11 +225,14 @@ def test_matmul():
     return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
 
   c = Tensor.custom_kernel(a, b, c, fxn=asm_kernel)[2]
-  ei = c.schedule()[0].lower()
+  linear = c.schedule_linear()
 
   ets = []
   with Context(DEBUG=2):
-    for _ in range(getenv("CNT", 5)): ets.append(ei.run(wait=True))
+    for _ in range(getenv("CNT", 5)):
+      start = GlobalCounters.time_sum_s
+      run_linear(linear)
+      ets.append(GlobalCounters.time_sum_s - start)
   print(f"REAL TFLOPS {N*N*N*2 / min(ets) * 1e-12:.2f}")
 
   if getenv("VERIFY", 1):

extra/gemm/simple_matmul.py

@@ -2,6 +2,7 @@ import numpy as np
 from tinygrad import dtypes, Tensor
 from tinygrad.helpers import getenv, get_single_element
 from tinygrad.dtype import _to_np_dtype
+from tinygrad.engine.realize import compile_linear
 from tinygrad.codegen.opt import OptOps
 
 dtype_in = (dtypes.half if getenv("HALF") else dtypes.bfloat16 if getenv("BFLOAT16") else
@@ -38,10 +39,10 @@ if __name__ == "__main__":
     c = a.matmul(b, dtype=acc_dtype).realize()
 
   if getenv("SHOULD_USE_TC"):
-    sched = a.matmul(b, dtype=acc_dtype).schedule()
-    ei = get_single_element(sched)
-    ei.lower()
-    assert any(opt.op is OptOps.TC for opt in ei.prg.p.applied_opts), f"TC not triggered, {ei.prg.p.applied_opts}"
+    linear = compile_linear(a.matmul(b, dtype=acc_dtype).schedule_linear())
+    call = get_single_element(list(linear.src))
+    applied_opts = call.src[0].src[0].arg.applied_opts
+    assert any(opt.op is OptOps.TC for opt in applied_opts), f"TC not triggered, {applied_opts}"
 
   ref = a.numpy().astype(np.float32) @ b.numpy().astype(np.float32)
   res = c.numpy()

extra/gemm/tinygrad_nv_matmul.py

@@ -1,7 +1,7 @@
-from tinygrad import Tensor, dtypes, Device
-from tinygrad.helpers import getenv, DEBUG
-from tinygrad.codegen.opt.kernel import Kernel, Opt, OptOps
-from tinygrad.engine.realize import CompiledRunner, ExecItem, get_program
+from tinygrad import Tensor, dtypes, Context
+from tinygrad.helpers import getenv
+from tinygrad.codegen.opt import Opt, OptOps
+from tinygrad.engine.realize import run_linear
 from dataclasses import replace
 
 N = 4096
@@ -11,9 +11,6 @@ if __name__ == "__main__":
   else:
     A, B = Tensor.empty(N, N, dtype=dtypes.float16), Tensor.empty(N, N, dtype=dtypes.float16)
   C = A.matmul(B)
-  si = C.schedule()[-1]
-  ast = si.ast
-  k = Kernel(ast, opts=Device[Device.DEFAULT].renderer)
   if getenv("GEMV"):
     opts = [
       Opt(op=OptOps.UNROLL, axis=0, amt=8),
@@ -28,10 +25,10 @@ if __name__ == "__main__":
       Opt(op=OptOps.LOCAL, axis=1, amt=2),
       Opt(op=OptOps.LOCAL, axis=0, amt=2),
     ]
-  k.apply_opts(opts)
-  prg = get_program(k.ast, k.opts, k.applied_opts)
-  new_src = prg.src
-  # can mod source here
-  prg = replace(prg, src=new_src)
-  ei = ExecItem(si.ast, [x.ensure_allocated() for x in si.bufs], si.metadata, prg=CompiledRunner(prg))
-  for i in range(5): ei.run(wait=True)
+  linear = C.schedule_linear()
+  call = linear.src[-1]
+  new_ast = call.src[0].replace(arg=replace(call.src[0].arg, opts_to_apply=tuple(opts)))
+  new_call = call.replace(src=(new_ast, *call.src[1:]))
+  linear = linear.replace(src=tuple(new_call if c is call else c for c in linear.src))
+  with Context(DEBUG=2):
+    for i in range(5): run_linear(linear)

extra/gemm/triton_nv_matmul.py

@@ -4,7 +4,9 @@ import triton.language as tl
 from triton.compiler import AttrsDescriptor, ASTSource, compile as triton_compile
 import numpy as np
 from tinygrad import Tensor, dtypes, Device
-from tinygrad.engine.realize import CompiledRunner, ExecItem, ProgramSpec
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
+from tinygrad.uop.ops import Ops, UOp, KernelInfo, ProgramInfo
 from tinygrad.helpers import getenv
 np.set_printoptions(suppress=True)
 
@@ -73,8 +75,11 @@ if __name__ == "__main__":
 
   A, B = Tensor.normal(M, K, std=1e-1, dtype=dtypes.float16).realize(), Tensor.normal(K, N, std=1e-1, dtype=dtypes.float16).realize()
   C = A.matmul(B)
-  sched = C.schedule()
-  si = sched[-1]
+  from tinygrad.uop.ops import Ops
+  linear, var_vals = C.linear_with_vars()
+  last_call = linear.src[-1]
+  ast = last_call.src[0]
+  bufs = [s.buffer for s in last_call.src[1:] if s.op is not Ops.BIND]
 
   src = compiled.asm["ptx"]
   # specify the shared memory here so we don't need to do it dynamically
@@ -85,22 +90,27 @@ if __name__ == "__main__":
   # remove debug sections
   src = src.split("\t.file")[0]
   assert '.extern .shared' not in src
-  prg = ProgramSpec("matmul_kernel", src, device=Device.DEFAULT,
-                global_size=[M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1], local_size=[32*compiled.metadata.num_warps, 1, 1],
-                mem_estimate=A.nbytes() + B.nbytes() + C.nbytes())
-  ei = ExecItem(si.ast, [x.ensure_allocated() for x in si.bufs], si.metadata, prg=CompiledRunner(prg))
+  info = ProgramInfo(name="matmul_kernel",
+                     global_size=(M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1), local_size=(32*compiled.metadata.num_warps, 1, 1))
+  sink = UOp.sink(arg=KernelInfo(name="matmul_kernel"))
+  prg_uop = to_program(UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR), UOp(Ops.SOURCE, arg=src)), arg=info),
+                       Device.default.renderer)
+  rt = get_runtime(Device.DEFAULT, prg_uop)
+  all_bufs = [x.ensure_allocated() for x in bufs]
+  prg_bufs = [all_bufs[i] for i in info.globals]
+  gsize, lsize = info.launch_dims({})
   tflops = []
   for i in range(5):
-    tm = ei.run(wait=True)
+    tm = rt(*[b._buf for b in prg_bufs], global_size=gsize, local_size=lsize, vals=info.vals({}), wait=True)
     tflops.append((2*M*K*N/tm)*1e-12)
   print(f"TFLOPS: {max(tflops):.2f}")
 
   # check correctness
   if getenv("VERIFY"):
-    from tinygrad.engine.realize import run_schedule
+    from tinygrad.engine.realize import run_linear
     triton_buf = np.frombuffer(si.bufs[0].as_memoryview(), np.float16).reshape(M,N)
     print(triton_buf)
-    run_schedule(sched)
+    run_linear(linear, var_vals)
     tinygrad_buf = np.frombuffer(si.bufs[0].as_memoryview(), np.float16).reshape(M,N)
     print(tinygrad_buf)
     np.testing.assert_allclose(triton_buf, tinygrad_buf)

extra/gemm/tvm_gemm.py

@@ -36,10 +36,10 @@ A = Tensor.rand(M, K, device="CPU")
 B = Tensor.rand(K, N, device="CPU")
 C = (A.reshape(M, 1, K) * B.permute(1,0).reshape(1, N, K)).sum(axis=2)
 
-sched = C.schedule()
+linear = C.schedule_linear()
 from tinygrad.codegen.opt.kernel import Kernel
 from tinygrad.device import CompilerOptions
-lin = Kernel(sched[-1].ast, CompilerOptions(has_local=False, supports_float4=False))
+lin = Kernel(linear.src[-1].src[0], CompilerOptions(has_local=False, supports_float4=False))
 lin.to_program()
 from tinygrad.runtime.ops_cpu import renderer
 src = renderer("mmult", lin.uops)

extra/llama_kernels/__init__.py

@@ -0,0 +1,51 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import Ops
+from tinygrad.runtime.support.compiler_amd import HIPCCCompiler
+
+FP8_MAX = 448.0
+NUM_WG, THREADS_PER_WG = 1024, 256
+
+# per-device abs max without allreduce
+@functools.cache
+def _local_abs_max_fxn(x_p, device):
+  x = Tensor(x_p, device=device)
+  inner = Tensor(x.uop.src[0]) if x.uop.op is Ops.MULTI else x
+  return (inner.abs().max(),)
+
+def local_abs_max(x:Tensor) -> Tensor:
+  param = x.as_param(0)
+  fxn = _local_abs_max_fxn(param.uop, x.device)
+  return Tensor(fxn[0].uop.call(x.uop).gettuple(0))
+
+def scalar_amax(amax_buf:Tensor) -> Tensor:
+  if isinstance(amax_buf.device, tuple):
+    return local_abs_max(amax_buf).detach()
+  return amax_buf.max().detach()
+
+def shard_shape(shape:tuple, axis:int, ndev:int) -> list:
+  s = list(shape)
+  s[axis] //= ndev
+  return s
+
+def dname_of(device) -> str:
+  if isinstance(device, tuple): return device[0].split(":")[0]
+  return device.split(":")[0] if isinstance(device, str) else device
+
+def alloc_like(shape, dtype, device, axis=None) -> Tensor:
+  if isinstance(device, tuple) and axis is not None:
+    return Tensor(Tensor.invalids(*shard_shape(shape, axis, len(device)), dtype=dtype, device=device).uop.multi(axis), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def alloc_local(shape, dtype, device, axis=None) -> Tensor:
+  if isinstance(device, tuple) and axis is not None:
+    return Tensor(Tensor.invalids(*shape, dtype=dtype, device=device).uop.multi(0), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def compile_hip(src:str, defines:list[str]):
+  return HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+
+def compile_cpp(cpp_dir:pathlib.Path, cpp_name:str, n_elems:int, hidden:int):
+  src = (cpp_dir/cpp_name).read_text()
+  return src, compile_hip(src, [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={hidden}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"])

extra/llama_kernels/cast_amax/__init__.py

@@ -0,0 +1,74 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, alloc_like, alloc_local, scalar_amax, dname_of
+
+# module-level mailbox: grad_xw13 UOp -> (grad_xw13_fp8 UOp, inv_scale UOp, new_amax UOp, store_effect)
+# lets cdna_asm_gemm's bwd reuse the fp8 companion produced by the fused silu_mul bwd kernel
+# instead of doing a redundant bf16 -> fp8 quantize.
+_grad_fp8_mailbox:dict = {}
+
+@functools.cache
+def _custom_fused_bwd_w13(grad_xw13:UOp, grad_xw13_fp8:UOp, grad_amax_buf:UOp,
+                          xw13:UOp, grad_x2:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
+  hidden = xw13.shape[2] // 2
+  n_elems = xw13.shape[0] * xw13.shape[1] * hidden
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 5 + n_elems * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(grad_xw13.base, grad_xw13_fp8.base, grad_amax_buf.base,
+                  xw13.base, grad_x2.base, amax_state.base, grad_amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_silu_mul_bwd_w13_{n_elems}", estimates=Estimates(ops=10*n_elems, mem=mem)))
+  src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_bwd_w13.cpp", n_elems, hidden)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+@functools.cache
+def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
+  # NOTE: grad_amax_state is plumbed through as an unused fwd input so the bwd kernel can read it via kernel.src
+  hidden = xw13.shape[2] // 2
+  n_elems = xw13.shape[0] * xw13.shape[1] * hidden
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 2 + n_elems + NUM_WG * 4
+  sink = UOp.sink(fp8_out.base, amax_buf.base, xw13.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_silu_mul_cast_amax_w13_{n_elems}", estimates=Estimates(ops=5*n_elems, mem=mem)))
+  src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_fwd_w13.cpp", n_elems, hidden)
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
+
+def _fused_quantize_bwd_w13(gradient:UOp, kernel:UOp):
+  _, _, xw13, amax_state, grad_amax_state = kernel.src[1:]
+  device = xw13.device
+  axis = xw13.axis if isinstance(device, tuple) else None
+  grad_xw13     = alloc_like(xw13.shape, dtypes.bfloat16, device, axis)
+  grad_xw13_fp8 = alloc_like(xw13.shape, dtypes.fp8e4m3,  device, axis)
+  grad_amax_buf = alloc_local((NUM_WG,), dtypes.float32,  device, axis)
+  grad_amax_state_t = Tensor(grad_amax_state, device=device)
+  fxn = functools.partial(_custom_fused_bwd_w13, dname=dname_of(device))
+  grad_xw13, grad_xw13_fp8, grad_amax_buf, *_ = Tensor.custom_kernel(
+    grad_xw13, grad_xw13_fp8, grad_amax_buf,
+    Tensor(xw13, device=device), Tensor(gradient, device=device).cast(dtypes.bfloat16),
+    Tensor(amax_state, device=device), grad_amax_state_t, fxn=fxn)
+  inv_scale = (grad_amax_state_t.float() + 1e-8) / FP8_MAX
+  new_grad_amax = scalar_amax(grad_amax_buf)
+  store_effect = grad_amax_state_t.uop.store(new_grad_amax.uop)
+  # Stash fp8 companion + amax store for cdna_asm_gemm's bwd to attach to grad_a.
+  _grad_fp8_mailbox[grad_xw13.uop] = (grad_xw13_fp8.uop, inv_scale.uop, new_grad_amax.uop, store_effect)
+  return (None, None, grad_xw13.uop, None, None)
+
+def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype, grad_amax_state:Tensor) -> tuple[Tensor, Tensor, Tensor]:
+  # NOTE: silu(xw1)*xw3 -> fp8 + amax over fused xw13 layout. Returns (fp8, inv_scale, new_amax)
+  # grad_amax_state: delayed amax for grad_xw13 fp8 quantization in the backward.
+  assert xw13.dtype == dtypes.bfloat16, f"expected bf16, got {xw13.dtype}"
+  MBS, SEQ, H2 = xw13.shape
+  assert H2 % 2 == 0, f"w13 last-axis must be even, got {H2}"
+  HIDDEN = H2 // 2
+  axis = xw13.uop.axis if isinstance(xw13.device, tuple) else None
+  fp8_out  = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,      xw13.device, axis)
+  amax_buf = alloc_local((NUM_WG,),         dtypes.float32, xw13.device, axis)
+  fxn = functools.partial(_custom_fused_cast_amax_w13, dname=dname_of(xw13.device))
+  fp8_out, amax_buf, *_ = Tensor.custom_kernel(fp8_out, amax_buf, xw13, amax_state, grad_amax_state,
+                                                fxn=fxn, grad_fxn=_fused_quantize_bwd_w13)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf)

extra/llama_kernels/cast_amax/cast_amax_bwd_w13.cpp

@@ -0,0 +1,98 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+#ifndef N_ELEMS
+#define N_ELEMS 234881024
+#endif
+#ifndef HIDDEN
+#define HIDDEN 14336
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC");
+
+// fused silu*mul backward, three outputs in a single HBM pass:
+//   1) bf16 grad_xw13      — consumed by downstream bf16 autograd chain
+//   2) fp8  grad_xw13_fp8  — delayed-scale quantize using grad_amax_state (mailbox to matmul bwd)
+//   3) fp32 grad_amax_buf  — per-WG partial |grad_xw13|, reduced into next step's grad_amax_state
+// grad_amax_state is read for the fp8 scale. The store of new_grad_amax into grad_amax_state's
+// buffer is built in Python as a separate effect and threaded into grad_a via .after(store).
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_silu_mul_bwd_w13(
+    __hip_bfloat16*       __restrict__ grad_xw13_out,        // bf16, 2*N_ELEMS
+    __hip_fp8_storage_t*  __restrict__ grad_xw13_fp8_out,    // fp8,  2*N_ELEMS
+    float*                __restrict__ grad_amax_buf,        // fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ xw13,                 // bf16, 2*N_ELEMS
+    const __hip_bfloat16* __restrict__ grad_x2,              // bf16, N_ELEMS
+    const float*          __restrict__ amax_state,           // fp32 scalar (fwd x2 amax)
+    const float*          __restrict__ grad_amax_state)      // fp32 scalar (delayed grad amax)
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float g_scale = FP8_MAX / (static_cast<float>(*grad_amax_state) + 1e-8f);
+  float local_max = 0.0f;
+
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    const int outer = base / HIDDEN;
+    const int inner = base % HIDDEN;
+    const int xw1_off = outer * 2 * HIDDEN + inner;
+    const int xw3_off = xw1_off + HIDDEN;
+
+    float4 x1_raw = *reinterpret_cast<const float4*>(&xw13[xw1_off]);
+    float4 x3_raw = *reinterpret_cast<const float4*>(&xw13[xw3_off]);
+    float4 g_raw  = *reinterpret_cast<const float4*>(&grad_x2[base]);
+
+    const __hip_bfloat16 *x1 = reinterpret_cast<const __hip_bfloat16*>(&x1_raw);
+    const __hip_bfloat16 *x3 = reinterpret_cast<const __hip_bfloat16*>(&x3_raw);
+    const __hip_bfloat16 *gv = reinterpret_cast<const __hip_bfloat16*>(&g_raw);
+
+    __hip_bfloat16 out1[VEC], out3[VEC];
+    __hip_fp8_storage_t fp8_1[VEC], fp8_3[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float f1 = static_cast<float>(x1[i]);
+      const float f3 = static_cast<float>(x3[i]);
+      const float fg = static_cast<float>(gv[i]);
+      const float sig = 1.0f / (1.0f + __expf(-f1));
+      const float silu = f1 * sig;
+      const float silu_prime = sig + silu * (1.0f - sig);
+      const float gs = fg * scale;
+      const float g1 = gs * silu_prime * f3;
+      const float g3 = gs * silu;
+      out1[i] = static_cast<__hip_bfloat16>(g1);
+      out3[i] = static_cast<__hip_bfloat16>(g3);
+      local_max = fmaxf(local_max, fmaxf(fabsf(g1), fabsf(g3)));
+      fp8_1[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g1 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
+      fp8_3[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g3 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
+    }
+
+    *reinterpret_cast<float4*>(&grad_xw13_out[xw1_off]) = *reinterpret_cast<float4*>(out1);
+    *reinterpret_cast<float4*>(&grad_xw13_out[xw3_off]) = *reinterpret_cast<float4*>(out3);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw1_off]) = *reinterpret_cast<uint64_t*>(fp8_1);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw3_off]) = *reinterpret_cast<uint64_t*>(fp8_3);
+  }
+
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) grad_amax_buf[wg] = sdata[0];
+}

extra/llama_kernels/cast_amax/cast_amax_fwd_w13.cpp

@@ -0,0 +1,79 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+#ifndef N_ELEMS
+#define N_ELEMS 234881024
+#endif
+#ifndef HIDDEN
+#define HIDDEN 14336
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC (so VEC loads don't straddle block boundary)");
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_silu_mul_cast_amax_w13(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG (per-WG amaxes)
+    const __hip_bfloat16* __restrict__ xw13,            // bf16, 2*N_ELEMS
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  float local_max = 0.0f;
+
+  // grid-stride over 8-element groups
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    // interleaved xw13 layout: xw1 and xw3 are not contiguous halves
+    const int outer = base / HIDDEN;
+    const int inner = base % HIDDEN;
+    const int xw1_off = outer * 2 * HIDDEN + inner;
+    const int xw3_off = xw1_off + HIDDEN;
+
+    float4 x1_raw = *reinterpret_cast<const float4*>(&xw13[xw1_off]);
+    float4 x3_raw = *reinterpret_cast<const float4*>(&xw13[xw3_off]);
+
+    const __hip_bfloat16 *x1 = reinterpret_cast<const __hip_bfloat16*>(&x1_raw);
+    const __hip_bfloat16 *x3 = reinterpret_cast<const __hip_bfloat16*>(&x3_raw);
+
+    __hip_fp8_storage_t out[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float f1 = static_cast<float>(x1[i]);
+      const float f3 = static_cast<float>(x3[i]);
+      const float silu = f1 / (1.0f + __expf(-f1));
+      const float x2 = silu * f3;
+      local_max = fmaxf(local_max, fabsf(x2));
+      const float x_scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, x2 * scale));
+      out[i] = __hip_cvt_float_to_fp8(x_scaled, __HIP_SATFINITE, __HIP_E4M3);
+    }
+
+    *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
+  }
+
+  // LDS tree reduction: per-workgroup amax
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+
+  if (tid == 0) amax_buf[wg] = sdata[0];
+}

extra/llama_kernels/fp8_transpose/__init__.py

@@ -0,0 +1,41 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import THREADS_PER_WG, alloc_like, dname_of, compile_hip
+
+TILE = 64
+
+@functools.cache
+def _custom_fp8_transpose(out:UOp, inp:UOp, dname:str) -> UOp:
+  M, N = inp.shape
+  num_wg = (M // TILE) * (N // TILE)
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(num_wg, "gidx0")
+  mem = M * N * 2  # one byte read + one byte write per element
+  sink = UOp.sink(out.base, inp.base, threads, workgroups,
+                  arg=KernelInfo(f"fp8_transpose_{M}_{N}",
+                                 estimates=Estimates(ops=M*N, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"fp8_transpose.cpp").read_text()
+  defines = [f"-DM_DIM={M}", f"-DN_DIM={N}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def fast_fp8_transpose(t:Tensor) -> Tensor:
+  assert t.ndim == 2, f"fast_fp8_transpose needs 2D input, got shape {t.shape}"
+  assert t.dtype in dtypes.fp8s, f"fast_fp8_transpose needs fp8 dtype, got {t.dtype}"
+  M, N = t.shape
+  assert M % TILE == 0 and N % TILE == 0, f"M={M}, N={N} must be multiples of {TILE}"
+
+  device = t.device
+  axis = t.uop.axis if isinstance(device, tuple) else None
+  out_axis = None
+  if axis == 0: out_axis = 1
+  elif axis == 1: out_axis = 0
+  elif axis is not None:
+    raise ValueError(f"fast_fp8_transpose: unsupported axis {axis}")
+
+  out = alloc_like((N, M), t.dtype, device, out_axis)
+  fxn = functools.partial(_custom_fp8_transpose, dname=dname_of(device))
+  out, _ = Tensor.custom_kernel(out, t, fxn=fxn)
+  return out