diff --git a/.github/workflows/benchmark.yml b/.github/workflows/benchmark.yml index a90d4c02b1..fb0e113284 100644 --- a/.github/workflows/benchmark.yml +++ b/.github/workflows/benchmark.yml @@ -81,85 +81,180 @@ jobs: # source /tmp/tinygrad_pytest_ci/bin/activate # pytest -nauto --durations=20 - testmacbenchmark: - name: Mac Benchmark - runs-on: [self-hosted, macOS] + sharedbenchmarks: + name: Benchmark (DEV=${{ matrix.dev }}) + runs-on: [self-hosted, "${{ matrix.dev == 'METAL' && 'macOS' || matrix.dev == 'AMD' && 'tinybox' || 'tinyboxgreen' }}"] + strategy: + fail-fast: false + matrix: + dev: ['METAL', 'AMD', 'NV'] timeout-minutes: 60 defaults: run: shell: bash -e -o pipefail {0} + env: + DEV: ${{ matrix.dev }} if: github.repository_owner == 'tinygrad' steps: - name: Checkout Code uses: actions/checkout@v6 - - name: Symlink models and datasets + - name: Setup (AMD) + if: ${{ matrix.dev == 'AMD' }} run: | - mkdir -p weights - mkdir -p extra/disassemblers - ln -s ~/tinygrad/extra/disassemblers/applegpu extra/disassemblers/applegpu - ln -s ~/tinygrad/weights/sd-v1-4.ckpt weights/sd-v1-4.ckpt - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz + ./extra/amdpci/setup_python_cap.sh + ./extra/hcq/hcq_smi.py amd rmmod + ./extra/hcq/hcq_smi.py amd kill_pids - name: setup staging db if: github.ref == 'refs/heads/update_benchmark_staging' run: | echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - name: reset process replay - run: python3.11 test/external/process_replay/reset.py - - name: Print macOS version - run: sw_vers - - name: Run Stable Diffusion - run: BENCHMARK_LOG=stable_diffusion JIT=1 ASSERT_MIN_STEP_TIME=720 python3.11 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing - - name: Run Stable Diffusion without fp16 - run: BENCHMARK_LOG=stable_diffusion_fp32 JIT=1 ASSERT_MIN_STEP_TIME=720 python3.11 examples/stable_diffusion.py --seed 0 --noshow --timing - - name: Run Stable Diffusion v2 - # TODO: very slow step time - run: BENCHMARK_LOG=stable_diffusion_v2 JIT=1 ASSERT_MIN_STEP_TIME=4500 python3.11 examples/sdv2.py --fp16 --seed 0 --noshow --timing - # process replay can't capture this, the graph is too large - - name: Run SDXL - run: BENCHMARK_LOG=stable_diffusion_xl ASSERT_MIN_STEP_TIME=5000 CAPTURE_PROCESS_REPLAY=0 JIT=1 python3.11 examples/sdxl.py --seed 0 --noshow --timing - - name: Run model inference benchmark - run: DEV=METAL NOCLANG=1 python3.11 test/external/external_model_benchmark.py - - name: Test speed vs torch - run: BIG=2 MPS=1 python3.11 test/speed/external_test_speed_v_torch.py + run: python3 test/external/process_replay/reset.py - name: Test tensor cores - run: DEV=METAL python3.11 test/opt/test_tensor_cores.py - - name: Run Tensor Core GEMM (float) - run: DEBUG=2 SHOULD_USE_TC=1 python3.11 extra/gemm/simple_matmul.py - - name: Run Tensor Core GEMM (half) - run: DEBUG=2 SHOULD_USE_TC=1 HALF=1 python3.11 extra/gemm/simple_matmul.py - - name: Run Tensor Core GEMM (bfloat16) - run: DEBUG=2 SHOULD_USE_TC=1 BFLOAT16=1 python3.11 extra/gemm/simple_matmul.py - - name: Fuzz Padded Tensor Core GEMM - run: DEV=METAL M_START=6 M_STOP=10 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=6 K_STOP=24 K_STEP=1 TC_OPT=2 DEBUG=2 python3.11 ./extra/gemm/fuzz_matmul.py + run: | + if [[ "${{ matrix.dev }}" == "METAL" ]]; then + python3 test/opt/test_tensor_cores.py + DEBUG=2 SHOULD_USE_TC=1 python3 extra/gemm/simple_matmul.py + DEBUG=2 SHOULD_USE_TC=1 HALF=1 python3 extra/gemm/simple_matmul.py + DEBUG=2 SHOULD_USE_TC=1 BFLOAT16=1 python3 extra/gemm/simple_matmul.py + M_START=6 M_STOP=10 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=6 K_STOP=24 K_STEP=1 TC_OPT=2 DEBUG=2 python3 ./extra/gemm/fuzz_matmul.py + elif [[ "${{ matrix.dev }}" == "NV" ]]; then + ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py + DEV=NV:PTX ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py + DEV=CUDA SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + DEV=CUDA SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + DEV=CUDA SHOULD_USE_TC=1 ALLOW_TF32=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py + DEV=CUDA SHOULD_USE_TC=1 FP8E4M3=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + DEV=NV:PTX SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + # TODO: too slow + # M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py + # DEV=NV:PTX M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py + else + python3 test/opt/test_tensor_cores.py + # TODO: this is flaky + # DEV=AMD:LLVM python3 test/opt/test_tensor_cores.py + SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + SHOULD_USE_TC=1 HALF=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py + # TODO: AMD compiler bug causes this to fail + # HSA=1 M_START=12 M_STOP=20 M_STEP=1 N_START=12 N_STOP=20 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 DEBUG=2 python3 ./extra/gemm/fuzz_matmul.py + fi + - name: Run model inference benchmark + # TODO: unstable on AMD + if: ${{ matrix.dev != 'AMD' }} + run: CAPTURE_PROCESS_REPLAY=0 NOCLANG=1 python3 test/external/external_model_benchmark.py + - name: Test speed vs torch + # TODO: unstable on AMD + if: ${{ matrix.dev != 'AMD' }} + env: + HALF: ${{ matrix.dev == 'NV' && '1' || '0' }} + run: CAPTURE_PROCESS_REPLAY=0 BIG=2 ${{ matrix.dev == 'METAL' && 'MPS=1' || 'TORCHCUDA=1' }} python3 test/speed/external_test_speed_v_torch.py + - name: Run Stable Diffusion + env: + ASSERT_MIN_STEP_TIME: ${{ matrix.dev == 'METAL' && '720' || matrix.dev == 'AMD' && '550' || '0' }} + run: BENCHMARK_LOG=stable_diffusion python3 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing + - name: Run SDXL + if: ${{ matrix.dev != 'NV' }} + env: + ASSERT_MIN_STEP_TIME: ${{ matrix.dev == 'METAL' && '5000' || matrix.dev == 'AMD' && '3200' || '2000' }} + run: BENCHMARK_LOG=stable_diffusion_xl CAPTURE_PROCESS_REPLAY=0 python3 examples/sdxl.py --seed 0 --noshow --timing - name: Run llama3.2 - run: BENCHMARK_LOG=llama32_3b-f16 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3.11 -m tinygrad.llm -m llama3.2:3b-f16 --benchmark --warmup + run: BENCHMARK_LOG=llama32_3b-f16 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 -m tinygrad.llm -m llama3.2:3b-f16 --benchmark --warmup + - name: Run qwen3.5 + # qwen3.5:35b-a3b doesn't fit on mac + if: ${{ matrix.dev != 'METAL' }} + run: BENCHMARK_LOG=qwen35_35b-a3b JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 -m tinygrad.llm -m qwen3.5:35b-a3b --benchmark --warmup - name: Run olmoe - run: BENCHMARK_LOG=olmoe JITBEAM=2 IGNORE_BEAM_CACHE=1 python3.11 -m tinygrad.llm -m olmoe --benchmark --warmup + # just metal for now + if: ${{ matrix.dev == 'METAL' }} + run: BENCHMARK_LOG=olmoe JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 -m tinygrad.llm -m olmoe --benchmark --warmup - name: Train MNIST - run: time PYTHONPATH=. TARGET_EVAL_ACC_PCT=96.0 python3.11 examples/beautiful_mnist.py - - # NOTE: this is failing in CI. it is not failing on my machine and I don't really have a way to debug it - # the error is "RuntimeError: Internal Error (0000000e:Internal Error)" - #- name: Run 10 CIFAR training steps - # run: BENCHMARK_LOG=cifar_10steps JIT=1 ASSERT_MIN_STEP_TIME=3000 STEPS=10 python3.11 examples/hlb_cifar10.py - #- name: Run 10 CIFAR training steps w HALF - # run: BENCHMARK_LOG=cifar_10steps_half JIT=2 ASSERT_MIN_STEP_TIME=3000 STEPS=10 DEFAULT_FLOAT=HALF python3.11 examples/hlb_cifar10.py - - #- name: Run 10 CIFAR training steps w BF16 - # run: STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3.11 examples/hlb_cifar10.py - # TODO: too slow - # - name: Run 10 CIFAR training steps w winograd - # run: BENCHMARK_LOG=cifar_10steps_wino JIT=1 ASSERT_MIN_STEP_TIME=150 WINO=1 STEPS=10 python3.11 examples/hlb_cifar10.py + run: time TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py + - name: Run 10 CIFAR training steps + env: + ASSERT_MIN_STEP_TIME: ${{ matrix.dev == 'NV' && '130' || matrix.dev == 'AMD' && '200' || '3000' }} + run: BENCHMARK_LOG=cifar_10steps STEPS=10 python3 examples/hlb_cifar10.py + - name: Run 10 CIFAR training steps w HALF + env: + ASSERT_MIN_STEP_TIME: ${{ matrix.dev == 'NV' && '120' || matrix.dev == 'AMD' && '230' || '3000' }} + run: BENCHMARK_LOG=cifar_10steps_half STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py - uses: actions/upload-artifact@v7 + if: ${{ matrix.dev != 'AMD' }} with: - name: Speed (Mac) + name: Speed (${{ matrix.dev }}) path: | onnx_inference_speed.csv - name: Run process replay tests uses: ./.github/actions/process-replay + tinyboxbenchmark: + name: Tinybox Benchmark (${{ matrix.dev }}) + runs-on: [self-hosted, Linux, "${{ matrix.dev == 'AMD' && 'tinybox' || 'tinyboxgreen' }}"] + strategy: + fail-fast: false + matrix: + dev: ['AMD', 'NV'] + timeout-minutes: 60 + defaults: + run: + shell: bash -e -o pipefail {0} + env: + DEV: ${{ matrix.dev }} + steps: + - name: Checkout Code + uses: actions/checkout@v6 + - name: Setup (AMD) + if: ${{ matrix.dev == 'AMD' }} + run: | + ./extra/amdpci/setup_python_cap.sh + ./extra/hcq/hcq_smi.py amd rmmod + ./extra/hcq/hcq_smi.py amd kill_pids + - name: Symlink models and datasets + run: | + mkdir -p weights + mkdir -p extra/datasets + ln -s /raid/weights/LLaMA-3 weights/LLaMA-3 + ln -s /raid/datasets/imagenet extra/datasets/imagenet + - name: setup staging db + if: github.ref == 'refs/heads/update_benchmark_staging' + run: | + echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV + rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal + - name: reset process replay + run: test/external/process_replay/reset.py + - name: Test tiny + run: | + DEBUG=2 python -m pytest -rA test/test_tiny.py + if [[ "${{ matrix.dev }}" == "NV" ]]; then + DEBUG=2 DEV=CUDA python -m pytest -rA test/test_tiny.py + fi + - name: Test speed vs theoretical + run: IGNORE_BEAM_CACHE=1 CCACHE=0 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20 + - name: Test benchmark allreduce + if: ${{ matrix.dev == 'NV' }} + run: python test/external/external_benchmark_multitensor_allreduce.py + - name: HEVC Decode Benchmark + if: ${{ matrix.dev == 'NV' }} + run: VALIDATE=1 MAX_FRAMES=100 ASSERT_FPS=1400 JITBEAM=1 PYTHONPATH=. python3 extra/hevc/decode.py + - name: Run LLaMA-3 8B on 4 GPUs with BEAM + run: BENCHMARK_LOG=llama3_beam_4gpu JITBEAM=2 IGNORE_BEAM_CACHE=1 CAPTURE_PROCESS_REPLAY=0 python3 examples/llama3.py --size 8B --shard 4 --model weights/LLaMA-3/8B-SF-DPO/ --benchmark --temperature 0 + - name: Run full CIFAR training w 1 GPU + run: time BENCHMARK_LOG=cifar DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py + - name: Run full CIFAR training steps w 6 GPUS + run: time BENCHMARK_LOG=cifar_6gpu CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py + - name: Run MLPerf resnet eval on training data + run: time BENCHMARK_LOG=resnet_eval MODEL=resnet python3 examples/mlperf/model_eval.py + - name: Run 10 MLPerf ResNet50 training steps (1 gpu) + run: BENCHMARK_LOG=resnet_10steps DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py + - name: Run 10 MLPerf ResNet50 training steps (6 gpu) + run: BENCHMARK_LOG=resnet_10steps_6gpu CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=1536 GPUS=6 MODEL=resnet python3 examples/mlperf/model_train.py + - name: Run 10 MLPerf Bert training steps (6 gpu) + # TODO: remove BERT_LAYERS once scheduler is fast + run: BENCHMARK_LOG=bert_10steps_6gpu 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 + uses: ./.github/actions/process-replay + testusbgpu: name: UsbGPU Benchmark runs-on: [self-hosted, macOS] @@ -196,337 +291,6 @@ jobs: - name: UsbGPU (USB4/TB) tiny tests run: PYTHONPATH=. DEV=PCI+NV:NAK python3.11 test/test_tiny.py - testnvidiabenchmark: - name: tinybox green Benchmark - runs-on: [self-hosted, Linux, tinyboxgreen] - timeout-minutes: 60 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Print nvidia-smi - run: nvidia-smi - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s /raid/weights/LLaMA-3 weights/LLaMA-3 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - - name: Run model inference benchmark - run: DEV=NV CAPTURE_PROCESS_REPLAY=0 NOCLANG=1 python3 test/external/external_model_benchmark.py - - name: Test speed vs torch - run: DEV=NV CAPTURE_PROCESS_REPLAY=0 HALF=1 BIG=2 TORCHCUDA=1 python3 test/speed/external_test_speed_v_torch.py - - name: Test speed vs theoretical - run: DEV=NV IGNORE_BEAM_CACHE=1 CCACHE=0 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20 - - name: Test benchmark allreduce - run: DEV=NV python test/external/external_benchmark_multitensor_allreduce.py - - name: Test tensor cores - run: | - DEV=NV ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py - DEV=NV:PTX ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py - - name: Run Tensor Core GEMM (CUDA) - run: | - DEV=CUDA SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - DEV=CUDA SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - DEV=CUDA SHOULD_USE_TC=1 ALLOW_TF32=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py - DEV=CUDA SHOULD_USE_TC=1 FP8E4M3=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - - name: Run Tensor Core GEMM (PTX) - run: DEV=NV:PTX SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - - name: Run Tensor Core GEMM (NV) - run: DEV=NV SHOULD_USE_TC=1 HALF=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - - name: Test DEV=NV - run: DEBUG=2 DEV=NV python -m pytest -rA test/test_tiny.py - - name: Test DEV=CUDA - run: DEBUG=2 DEV=CUDA python -m pytest -rA test/test_tiny.py - - name: Run Stable Diffusion - run: BENCHMARK_LOG=stable_diffusion DEV=NV python3 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing - # TODO: too slow - # - name: Run SDXL - # run: BENCHMARK_LOG=stable_diffusion_xl ASSERT_MIN_STEP_TIME=2000 CAPTURE_PROCESS_REPLAY=0 DEV=NV CAPTURE_PROCESS_REPLAY=0 python3 examples/sdxl.py --seed 0 --noshow --timing - - name: Run llama3.2 - run: DEV=NV BENCHMARK_LOG=llama32_3b-f16 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 -m tinygrad.llm -m llama3.2:3b-f16 --benchmark --warmup - - name: Run qwen3.5 - run: DEV=NV BENCHMARK_LOG=qwen35_35b-a3b JITBEAM=2 IGNORE_BEAM_CACHE=1 CAPTURE_PROCESS_REPLAY=0 python3 -m tinygrad.llm -m qwen3.5:35b-a3b --benchmark --warmup - - name: Run LLaMA-3 8B on 4 GPUs with BEAM - run: BENCHMARK_LOG=llama3_beam_4gpu DEV=NV JITBEAM=2 IGNORE_BEAM_CACHE=1 CAPTURE_PROCESS_REPLAY=0 python3 examples/llama3.py --size 8B --shard 4 --model weights/LLaMA-3/8B-SF-DPO/ --benchmark --temperature 0 - # - name: Run LLaMA-3 8B on 6 GPUs - # run: DEV=NV CAPTURE_PROCESS_REPLAY=0 python3 examples/llama3.py --size 8B --shard 6 --model weights/LLaMA-3/8B-SF-DPO/ --benchmark --temperature 0 - # - name: Run LLaMA-2 70B - # run: DEV=NV CAPTURE_PROCESS_REPLAY=0 MAX_CONTEXT=256 python3 examples/llama.py --gen 2 --size 70B --shard 6 --prompt "Hello." --count 10 --temperature 0 --timing - - uses: actions/upload-artifact@v7 - with: - name: Speed (NVIDIA) - path: | - onnx_inference_speed.csv - - name: Run process replay tests - uses: ./.github/actions/process-replay - - testmorenvidiabenchmark: - name: tinybox green Training Benchmark - runs-on: [self-hosted, Linux, tinyboxgreen] - timeout-minutes: 60 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/weights/LLaMA weights/LLaMA - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/mixtral-8x7b-32kseqlen weights/mixtral-8x7b-32kseqlen - ln -s /raid/weights/LLaMA-2 weights/LLaMA-2 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - # TODO: too slow - # - name: Fuzz Padded Tensor Core GEMM (NV) - # run: DEV=NV M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py - # TODO: too slow - # - name: Fuzz Padded Tensor Core GEMM (PTX) - # run: DEV=NV:PTX M_START=12 M_STOP=20 M_STEP=1 N_START=6 N_STOP=10 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 python3 ./extra/gemm/fuzz_matmul.py - - name: HEVC Decode Benchmark - run: VALIDATE=1 MAX_FRAMES=100 ASSERT_FPS=1400 JITBEAM=1 DEV=NV PYTHONPATH=. python3 extra/hevc/decode.py - - name: Train MNIST - run: time PYTHONPATH=. DEV=NV TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py - - name: Run 10 CIFAR training steps - run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=130 DEV=NV STEPS=10 python3 examples/hlb_cifar10.py - - name: Run 10 CIFAR training steps w HALF - run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=120 DEV=NV STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py - - name: Run 10 CIFAR training steps w BF16 - run: BENCHMARK_LOG=cifar_10steps_bf16 ASSERT_MIN_STEP_TIME=120 DEV=NV STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3 examples/hlb_cifar10.py - # - name: Run 10 CIFAR training steps w winograd - # run: BENCHMARK_LOG=cifar_10steps_half_wino ASSERT_MIN_STEP_TIME=350 DEV=NV WINO=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py - - name: Run full CIFAR training w 1 GPU - run: time BENCHMARK_LOG=cifar DEV=NV DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - - name: Run full CIFAR training steps w 6 GPUS - run: time BENCHMARK_LOG=cifar_6gpu CAPTURE_PROCESS_REPLAY=0 DEV=NV DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - - name: Run MLPerf resnet eval on training data - run: time BENCHMARK_LOG=resnet_eval DEV=NV MODEL=resnet python3 examples/mlperf/model_eval.py - - name: Run 10 MLPerf ResNet50 training steps (1 gpu) - run: BENCHMARK_LOG=resnet_10steps DEV=NV DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py - - name: Run 10 MLPerf ResNet50 training steps (6 gpu) - run: BENCHMARK_LOG=resnet_10steps_6gpu DEV=NV CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=1536 GPUS=6 MODEL=resnet python3 examples/mlperf/model_train.py - - name: Run 10 MLPerf Bert training steps (6 gpu) - # 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 - uses: ./.github/actions/process-replay - - testamdbenchmark: - name: tinybox red Benchmark - runs-on: [self-hosted, Linux, tinybox] - timeout-minutes: 60 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Setcap to python - run: ./extra/amdpci/setup_python_cap.sh - - name: Remove amd modules - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd rmmod - - name: Kill stale pids - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd kill_pids - #- name: Insert amdgpu - # run: sudo modprobe amdgpu - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/LLaMA-3 weights/LLaMA-3 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - #- name: setup perflevel - # run: | - # examples/mlperf/training_submission_v4.1/tinycorp/benchmarks/bert/implementations/tinybox_red/setup.sh - # rocm-smi - #- name: Show off tinybox - # run: /opt/rocm/bin/rocm-bandwidth-test - # TODO: unstable on AMD - #- name: Run model inference benchmark - # run: LD_PRELOAD="/opt/rocm/lib/libhsa-runtime64.so" HSA=1 NOCLANG=1 python3 test/external/external_model_benchmark.py - # TODO: unstable on AMD - #- name: Test speed vs torch - # run: | - # python3 -c "import torch; print(torch.__version__)" - # LD_PRELOAD="/opt/rocm/lib/libhsa-runtime64.so" HSA=1 BIG=2 TORCHCUDA=1 python3 test/speed/external_test_speed_v_torch.py - - name: Test speed vs theoretical - run: DEV=AMD IGNORE_BEAM_CACHE=1 CCACHE=0 BEAM_DEBUG=1 DEBUG=1 python -m pytest -rA test/external/speed_v_theoretical.py --durations=20 - - name: Test tensor cores (no LLVM) - run: DEV=AMD python3 test/opt/test_tensor_cores.py - # TODO: this is flaky - # - name: Test tensor cores AMD:LLVM - # run: DEV=AMD:LLVM python3 test/opt/test_tensor_cores.py - - name: Run Tensor Core GEMM (AMD) - run: | - DEV=AMD SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - DEV=AMD SHOULD_USE_TC=1 HALF=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py - - name: Test DEV=AMD - run: DEBUG=2 DEV=AMD python -m pytest -rA test/test_tiny.py - #- name: Test HIP=1 - # run: DEBUG=2 HIP=1 python -m pytest -rA test/test_tiny.py - # TODO: AMD compiler bug causes this to fail - #- name: Fuzz Padded Tensor Core GEMM - # run: HSA=1 M_START=12 M_STOP=20 M_STEP=1 N_START=12 N_STOP=20 N_STEP=1 K_START=28 K_STOP=36 K_STEP=1 HALF=1 TC_OPT=2 DEBUG=2 python3 ./extra/gemm/fuzz_matmul.py - #- name: Remove amdgpu - # run: sleep 10 && sudo rmmod amdgpu # sleep a bit to let the driver unload the prev pid. - - name: Test AM cold start time - run: time DEV=AMD AM_RESET=1 python3 test/test_tiny.py TestTiny.test_plus - - name: Test AM warm start time - run: time DEV=AMD python3 test/test_tiny.py TestTiny.test_plus - - name: Run Stable Diffusion - run: BENCHMARK_LOG=stable_diffusion ASSERT_MIN_STEP_TIME=550 DEV=AMD python3 examples/stable_diffusion.py --fp16 --seed 0 --noshow --timing - - name: Run SDXL - run: BENCHMARK_LOG=stable_diffusion_xl ASSERT_MIN_STEP_TIME=3200 CAPTURE_PROCESS_REPLAY=0 DEV=AMD python3 examples/sdxl.py --seed 0 --noshow --timing - - name: Run llama3.2 - run: DEV=AMD BENCHMARK_LOG=llama32_3b-f16 JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 -m tinygrad.llm -m llama3.2:3b-f16 --benchmark --warmup - - name: Run qwen3.5 - run: DEV=AMD BENCHMARK_LOG=qwen35_35b-a3b JITBEAM=2 IGNORE_BEAM_CACHE=1 CAPTURE_PROCESS_REPLAY=0 python3 -m tinygrad.llm -m qwen3.5:35b-a3b --benchmark --warmup - - name: Run LLaMA-3 8B on 4 GPUs with BEAM - run: BENCHMARK_LOG=llama3_beam_4gpu DEV=AMD JITBEAM=2 IGNORE_BEAM_CACHE=1 CAPTURE_PROCESS_REPLAY=0 python3 examples/llama3.py --size 8B --shard 4 --model weights/LLaMA-3/8B-SF-DPO/ --benchmark --temperature 0 - # - name: Run LLaMA-3 8B on 6 GPUs - # run: DEV=AMD CAPTURE_PROCESS_REPLAY=0 python3 examples/llama3.py --size 8B --shard 6 --model weights/LLaMA-3/8B-SF-DPO/ --benchmark --temperature 0 - #- name: Restore amdgpu - # run: sudo modprobe amdgpu - # - name: Run LLaMA-2 70B - # run: DEV=AMD CAPTURE_PROCESS_REPLAY=0 python3 examples/llama.py --gen 2 --size 70B --shard 6 --prompt "Hello." --count 10 --temperature 0 --timing - - name: Run process replay tests - uses: ./.github/actions/process-replay - - testmoreamdbenchmark: - name: tinybox red Training Benchmark - runs-on: [self-hosted, Linux, tinybox] - timeout-minutes: 60 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Setcap to python - run: ./extra/amdpci/setup_python_cap.sh - - name: Remove amd modules - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd rmmod - - name: Kill stale pids - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd kill_pids - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/weights/LLaMA weights/LLaMA - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/mixtral-8x7b-32kseqlen weights/mixtral-8x7b-32kseqlen - ln -s /raid/weights/LLaMA-2 weights/LLaMA-2 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - - name: Test GPU crash recovery - run: DEV=AMD python3 -m pytest -rA test/external/external_test_gpu_crash.py - - name: Train MNIST - run: time PYTHONPATH=. DEV=AMD TARGET_EVAL_ACC_PCT=96.0 python3 examples/beautiful_mnist.py - - name: Run 10 CIFAR training steps - run: BENCHMARK_LOG=cifar_10steps ASSERT_MIN_STEP_TIME=200 DEV=AMD STEPS=10 python3 examples/hlb_cifar10.py - - name: Run 10 CIFAR training steps w HALF - run: BENCHMARK_LOG=cifar_10steps_half ASSERT_MIN_STEP_TIME=230 DEV=AMD STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py - # - name: Run 10 CIFAR training steps w BF16 - # run: BENCHMARK_LOG=cifar_10steps_bf16 ASSERT_MIN_STEP_TIME=288 DEV=AMD STEPS=10 DEFAULT_FLOAT=BFLOAT16 python3 examples/hlb_cifar10.py - # TODO: too slow - # - name: Run 10 CIFAR training steps w winograd - # run: BENCHMARK_LOG=cifar_10steps_half_wino ASSERT_MIN_STEP_TIME=66 DEV=AMD WINO=1 STEPS=10 DEFAULT_FLOAT=HALF python3 examples/hlb_cifar10.py - - name: Run full CIFAR training w 1 GPU - run: time BENCHMARK_LOG=cifar DEV=AMD DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - - name: Run full CIFAR training steps w 6 GPUS - run: time BENCHMARK_LOG=cifar_6gpu DEV=AMD DEFAULT_FLOAT=HALF STEPS=350 BS=1536 GPUS=6 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - # TODO: broken on some of the machines - #- 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 - uses: ./.github/actions/process-replay - - testmlperfamdbenchmark: - name: tinybox red MLPerf Benchmark - runs-on: [self-hosted, Linux, tinybox] - timeout-minutes: 60 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Setcap to python - run: ./extra/amdpci/setup_python_cap.sh - - name: Remove amd modules - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd rmmod - - name: Kill stale pids - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd kill_pids - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/weights/LLaMA weights/LLaMA - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/mixtral-8x7b-32kseqlen weights/mixtral-8x7b-32kseqlen - ln -s /raid/weights/LLaMA-2 weights/LLaMA-2 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - - name: Run MLPerf resnet eval - run: time BENCHMARK_LOG=resnet_eval DEV=AMD MODEL=resnet python3 examples/mlperf/model_eval.py - - name: Run 10 MLPerf ResNet50 training steps (1 gpu) - run: BENCHMARK_LOG=resnet_10steps DEV=AMD DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py - - name: Run 10 MLPerf ResNet50 training steps (6 gpu) - run: BENCHMARK_LOG=resnet_10steps_6gpu DEV=AMD CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=1536 GPUS=6 MODEL=resnet python3 examples/mlperf/model_train.py - - name: Run 10 MLPerf Bert training steps (6 gpu) - # 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 - uses: ./.github/actions/process-replay - testcommalatest: name: comma Benchmark (0.11.0) runs-on: [self-hosted, Linux, comma] @@ -648,31 +412,28 @@ jobs: - 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 + driverbenchmarks: + name: PCI Driver Benchmark (DEV=${{ matrix.dev }}) runs-on: [self-hosted, Linux, tinyboxrandom] + strategy: + fail-fast: false + matrix: + dev: ['AMD', 'NV'] timeout-minutes: 20 defaults: run: shell: bash -e -o pipefail {0} + env: + DEV: ${{ matrix.dev }} if: github.repository_owner == 'tinygrad' steps: - name: Checkout Code uses: actions/checkout@v6 - - name: Setcap to python - run: ./extra/amdpci/setup_python_cap.sh - - name: Remove amd modules - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd rmmod - - name: Kill stale pids - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py amd kill_pids - - name: Symlink models and datasets + - name: Setup run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/weights/LLaMA weights/LLaMA - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/mixtral-8x7b-32kseqlen weights/mixtral-8x7b-32kseqlen - ln -s /raid/weights/LLaMA-2 weights/LLaMA-2 + ./extra/amdpci/setup_python_cap.sh + ./extra/hcq/hcq_smi.py ${{ matrix.dev == 'AMD' && 'amd' || 'nv' }} rmmod + ./extra/hcq/hcq_smi.py ${{ matrix.dev == 'AMD' && 'amd' || 'nv' }} kill_pids mkdir -p extra/datasets ln -s /raid/datasets/imagenet extra/datasets/imagenet - name: setup staging db @@ -683,104 +444,47 @@ jobs: - name: reset process replay run: test/external/process_replay/reset.py - name: Test driver cold start time - run: time DEBUG=3 DEV=AMD AM_RESET=1 python3 test/test_tiny.py TestTiny.test_plus + run: time DEBUG=3 AM_RESET=1 python3 test/test_tiny.py TestTiny.test_plus - name: Test driver warm start time - run: time DEBUG=3 DEV=AMD python3 test/test_tiny.py TestTiny.test_plus + if: ${{ matrix.dev == 'AMD' }} + run: time DEBUG=3 python3 test/test_tiny.py TestTiny.test_plus - name: Test GPU crash recovery - run: DEV=AMD python3 -m pytest -rA test/external/external_test_gpu_crash.py - # Fails on 9070 - # - name: Test tensor cores - # run: | - # DEV=AMD python3 test/test_linearizer.py test/opt/test_tensor_cores.py - # DEV=AMD:LLVM python3 test/test_linearizer.py test/opt/test_tensor_cores.py - # DEV=AMD SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py - - name: Run Tensor Core GEMM (AMD) - run: DEV=AMD SHOULD_USE_TC=1 HALF=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py - - name: Test DEV=AMD - run: DEBUG=2 DEV=AMD python -m pytest -rA test/test_tiny.py + if: ${{ matrix.dev == 'AMD' }} + run: python3 -m pytest -rA test/external/external_test_gpu_crash.py + - name: Test tensor cores + run: | + if [[ "${{ matrix.dev }}" == "AMD" ]]; then + # Fails on 9070 + # python3 test/test_linearizer.py test/opt/test_tensor_cores.py + # DEV=AMD:LLVM python3 test/test_linearizer.py test/opt/test_tensor_cores.py + # SHOULD_USE_TC=1 BFLOAT16=1 DEBUG=2 python3 extra/gemm/simple_matmul.py + SHOULD_USE_TC=1 HALF=1 DEBUG=2 ATOL=2e-2 python3 extra/gemm/simple_matmul.py + else + ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py + fi - name: Test DISK copy time - run: DEV=AMD TESTFILE=/raid/downloads/llama3-8b-sfr/model-00001-of-00004.safetensors python3 test/external/external_benchmark_disk_raw.py + run: TESTFILE=/raid/downloads/llama3-8b-sfr/model-00001-of-00004.safetensors python3 test/external/external_benchmark_disk_raw.py - name: Test CPU copy time run: | - DEV=AMD GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyDefaulttoCPUJit - DEV=AMD GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyCPUtoDefaultJit + GRAPH_ONE_KERNEL=1 NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyDefaulttoCPUJit + GRAPH_ONE_KERNEL=1 NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyCPUtoDefaultJit - name: Run full CIFAR training w 1 GPU - run: time BENCHMARK_LOG=cifar DEV=AMD DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - # - name: Run 10 MLPerf ResNet50 training steps (1 gpu) - # run: BENCHMARK_LOG=resnet_10steps DEV=AMD MNISTMOCK=1 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py + run: time BENCHMARK_LOG=cifar DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py + - name: Run 10 MLPerf ResNet50 training steps (1 gpu) + if: ${{ matrix.dev == 'NV' }} + run: BENCHMARK_LOG=resnet_10steps MNISTMOCK=1 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py - name: Run 10 MLPerf Bert training steps (1 gpu) # TODO: remove BERT_LAYERS once scheduler is fast - run: BENCHMARK_LOG=bert_10steps DEV=AMD CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=66 GPUS=1 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py + run: BENCHMARK_LOG=bert_10steps CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=66 GPUS=1 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py - name: Remote run: | pkill -f 'extra/remote/serve.py' || true PYTHONPATH=. python3 extra/remote/serve.py 6482 & sleep 1 - DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 DEV=PCI+AMD python3 test/test_tiny.py - 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 - uses: ./.github/actions/process-replay - - testgreendriverbenchmark: - name: NV Benchmark - runs-on: [self-hosted, Linux, tinyboxrandom] - timeout-minutes: 20 - defaults: - run: - shell: bash -e -o pipefail {0} - if: github.repository_owner == 'tinygrad' - steps: - - name: Checkout Code - uses: actions/checkout@v6 - - name: Setcap to python - run: ./extra/amdpci/setup_python_cap.sh - - name: Remove nv modules - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py nv rmmod - - name: Kill stale pids - run: PYTHONPATH=. ./extra/hcq/hcq_smi.py nv kill_pids - - name: Symlink models and datasets - run: | - mkdir -p weights - ln -s ~/tinygrad/weights/bpe_simple_vocab_16e6.txt.gz weights/bpe_simple_vocab_16e6.txt.gz - ln -s ~/tinygrad/weights/LLaMA weights/LLaMA - ln -s ~/tinygrad/extra/datasets/cifar-10-python.tar.gz extra/datasets/cifar-10-python.tar.gz - ln -s /raid/weights/mixtral-8x7b-32kseqlen weights/mixtral-8x7b-32kseqlen - ln -s /raid/weights/LLaMA-2 weights/LLaMA-2 - mkdir -p extra/datasets - ln -s /raid/datasets/imagenet extra/datasets/imagenet - - name: setup staging db - if: github.ref == 'refs/heads/update_benchmark_staging' - run: | - echo "CACHEDB=/tmp/staging.db" >> $GITHUB_ENV - rm -f /tmp/staging.db /tmp/staging.db-shm /tmp/staging.db-wal - - name: reset process replay - run: test/external/process_replay/reset.py - - name: Test driver start time - run: time DEBUG=3 DEV=NV python3 test/test_tiny.py TestTiny.test_plus - - name: Test tensor cores - run: DEV=NV ALLOW_TF32=1 python3 test/opt/test_tensor_cores.py - - name: Test DISK copy time - run: DEV=NV TESTFILE=/raid/downloads/llama3-8b-sfr/model-00001-of-00004.safetensors python3 test/external/external_benchmark_disk_raw.py - - name: Test CPU copy time - run: | - DEV=NV GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyDefaulttoCPUJit - DEV=NV GRAPH_ONE_KERNEL=1 PYTHONPATH=. NSZ=8192 python3 test/speed/external_test_copy_speed.py TestCopySpeed.testCopyCPUtoDefaultJit - - name: Test LLAMA-3 - run: BENCHMARK_LOG=llama3_beam DEV=NV JITBEAM=2 IGNORE_BEAM_CACHE=1 python3 examples/llama3.py --size 8B --benchmark --temperature 0 - - name: Run full CIFAR training w 1 GPU - run: time BENCHMARK_LOG=cifar DEV=NV DEFAULT_FLOAT=HALF STEPS=1000 TARGET_EVAL_ACC_PCT=93.0 python3 examples/hlb_cifar10.py - - name: Run 10 MLPerf ResNet50 training steps (1 gpu) - run: BENCHMARK_LOG=resnet_10steps DEV=NV MNISTMOCK=1 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=256 GPUS=1 MODEL=resnet python3 examples/mlperf/model_train.py - - name: Run 10 MLPerf Bert training steps (1 gpu) - # TODO: remove BERT_LAYERS once scheduler is fast - run: BENCHMARK_LOG=bert_10steps DEV=NV CAPTURE_PROCESS_REPLAY=0 DEFAULT_FLOAT=HALF BENCHMARK=10 BS=66 GPUS=1 BERT_LAYERS=2 MODEL=bert python3 examples/mlperf/model_train.py - - name: Remote - run: | - pkill -f 'extra/remote/serve.py' || true - PYTHONPATH=. python3 extra/remote/serve.py 6483 & - sleep 1 - DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6483 DEV=NV python3 test/test_tiny.py + DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 python3 test/test_tiny.py + if [[ "${{ matrix.dev }}" == "AMD" ]]; then + DEBUG=2 PYTHONPATH=. REMOTE=127.0.0.1:6482 AM_RESET=1 AMD_AQL=1 python3 test/test_tiny.py + fi pkill -f 'extra/remote/serve.py' || true - name: Run process replay tests uses: ./.github/actions/process-replay diff --git a/examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/profile.sh b/examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/profile.sh index b52fed8a9f..3f78ebe7c4 100755 --- a/examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/profile.sh +++ b/examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/profile.sh @@ -3,4 +3,4 @@ export BENCHMARK=${BENCHMARK:-5} export EVAL_BS=0 VIZ=${VIZ:--1} FULL_LAYERS=1 DEBUG=${DEBUG:--0} examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama31_8b/implementations/tinybox_8xMI350X/dev_beam.sh SRC="AMD"; [[ $DEV == NULL* ]] && SRC="NULL" -[ $BENCHMARK -gt 3 ] && python -m tinygrad.viz.cli -s "$SRC" -t --interval "train @ 2" "train @ 3" +[ "$BENCHMARK" -le 3 ] || python -m tinygrad.viz.cli -s "$SRC" -t --interval "train @ 2" "train @ 3" diff --git a/extra/gemm/cdna_asm_gemm.py b/extra/gemm/cdna_asm_gemm.py index cb47ce8769..2fa695b005 100644 --- a/extra/gemm/cdna_asm_gemm.py +++ b/extra/gemm/cdna_asm_gemm.py @@ -2899,7 +2899,7 @@ def custom_mx_gemm_bw(gradient:UOp, kernel:UOp, has_w_post:bool, w_stored:bool=F g = Tensor(gradient, device=aq.device)[:aq.shape[0]].reshape(aq.shape[0]*aq.shape[1], bq.shape[0]).cast(dtypes.bfloat16) grad_a = asm_gemm(g, b_phys, mx=True) - grad_b = asm_gemm(g.T, a_phys, mx=True) + grad_b = asm_gemm(g.T, a_phys, mx=True, a_pretranspose=g) grad_a = (grad_a * _mx_block_scale(ae8)).reshape(aq.shape) if not w_stored: grad_b = grad_b * _mx_block_scale(be8) @@ -2909,7 +2909,8 @@ def custom_mx_gemm_bw(gradient:UOp, kernel:UOp, has_w_post:bool, w_stored:bool=F # ** main gemm function def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=None, grad_amax_state:Tensor|None=None, - w_post_scale:Tensor|None=None, mx:bool=False, mx_scales:tuple|None=None, mx_w_stored:bool=False, g_scale:Tensor|None=None) -> Tensor: + w_post_scale:Tensor|None=None, mx:bool=False, mx_scales:tuple|None=None, mx_w_stored:bool=False, g_scale:Tensor|None=None, + a_pretranspose: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 @@ -2946,6 +2947,11 @@ def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=N if mx_scales is not None: a_si, a_e8, b_si, b_e8 = mx_scales a_q, b_q = a.reshape(-1, a.shape[-1]), b.T + elif (a_pretranspose is not None and getenv("FUSED_GRAD_QUANTIZE", 0) and a_pretranspose.dtype == dtypes.bfloat16 + and a_pretranspose.shape[0] % 32 == 0 and a_pretranspose.shape[1] % 256 == 0): + from extra.llama_kernels.transpose_quantize_mxfp8 import transpose_quantize_mxfp8 + a_q, a_e8, a_si = transpose_quantize_mxfp8(a_pretranspose) + b_q, b_e8, b_si = quantize_mxfp8(b.T) else: a_q, a_e8, a_si = quantize_mxfp8(a.reshape(-1, a.shape[-1])) b_q, b_e8, b_si = quantize_mxfp8(b.T) diff --git a/extra/llama_kernels/transpose_quantize_mxfp8/__init__.py b/extra/llama_kernels/transpose_quantize_mxfp8/__init__.py new file mode 100644 index 0000000000..e5b43bebc8 --- /dev/null +++ b/extra/llama_kernels/transpose_quantize_mxfp8/__init__.py @@ -0,0 +1,37 @@ +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_N = THREADS_PER_WG # 256 +BLK = 32 + +@functools.cache +def _custom_transpose_quantize_mxfp8(q:UOp, e8:UOp, g:UOp, dname:str) -> UOp: + M, N = g.shape + num_wg = (M // BLK) * (N // TILE_N) + threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(num_wg, "gidx0") + mem = M * N * 2 + M * N + (M // BLK) * N # read bf16, write fp8 + e8 + sink = UOp.sink(q.base, e8.base, g.base, threads, workgroups, + arg=KernelInfo(f"transpose_quantize_mxfp8_{M}_{N}", estimates=Estimates(ops=M*N, mem=mem))) + src = (pathlib.Path(__file__).parent/"transpose_quantize_mxfp8.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 transpose_quantize_mxfp8(g:Tensor) -> tuple[Tensor, Tensor, Tensor]: + # fused g.T quantize: returns (q, e8, si) == quantize_mxfp8(g.T) — q (N,M) fp8, e8 (N, M/32), si packed (M/128, N) + assert g.ndim == 2 and g.dtype == dtypes.bfloat16, f"{g.shape} {g.dtype}" + from extra.gemm.cdna_asm_gemm import FP8_DTYPE, mx_pack + M, N = g.shape + assert M % BLK == 0 and N % TILE_N == 0, f"M={M} must%{BLK}, N={N} must%{TILE_N}" + device = g.device + axis = g.uop.axis if isinstance(device, tuple) else None + out_axis = None if axis is None else (1 if axis == 0 else 0) + q = alloc_like((N, M), FP8_DTYPE, device, out_axis) + e8 = alloc_like((N, M // BLK), dtypes.uint8, device, out_axis) + fxn = functools.partial(_custom_transpose_quantize_mxfp8, dname=dname_of(device)) + q, e8, *_ = Tensor.custom_kernel(q, e8, g, fxn=fxn) + return q, e8, mx_pack(e8) diff --git a/extra/llama_kernels/transpose_quantize_mxfp8/transpose_quantize_mxfp8.cpp b/extra/llama_kernels/transpose_quantize_mxfp8/transpose_quantize_mxfp8.cpp new file mode 100644 index 0000000000..157e57e0ba --- /dev/null +++ b/extra/llama_kernels/transpose_quantize_mxfp8/transpose_quantize_mxfp8.cpp @@ -0,0 +1,62 @@ +#include +#include +#include + +#ifndef M_DIM +#define M_DIM 8192 +#endif +#ifndef N_DIM +#define N_DIM 14336 +#endif +#ifndef THREADS_PER_WG +#define THREADS_PER_WG 256 +#endif + +constexpr int BLK = 32; +constexpr int TILE_M = BLK; // one mxfp8 block along M per tile +constexpr int TILE_N = THREADS_PER_WG; // 256, one output column per thread +constexpr int LDS_STRIDE = TILE_N + 1; // +1 pad: stride 257 ≡ 1 (mod 32) -> conflict-free column reads +constexpr int N_TILES_N = N_DIM / TILE_N; +constexpr float FP8_MAX = 448.0f; + +static_assert(M_DIM % TILE_M == 0, "M_DIM must be a multiple of 32"); +static_assert(N_DIM % TILE_N == 0, "N_DIM must be a multiple of 256"); + +extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void +transpose_quantize_mxfp8(__hip_fp8_storage_t* __restrict__ q, // (N_DIM, M_DIM) + uint8_t* __restrict__ e8_out, // (N_DIM, M_DIM/32) + const __hip_bfloat16* __restrict__ g) // (M_DIM, N_DIM) +{ + __shared__ __hip_bfloat16 lds[TILE_M * LDS_STRIDE]; + const int tid = threadIdx.x; + const int tile_m = blockIdx.x / N_TILES_N; // which 32-block along M + const int tile_n = blockIdx.x % N_TILES_N; + + #pragma unroll + for (int mm = 0; mm < TILE_M; mm++) + lds[mm * LDS_STRIDE + tid] = g[(long long)(tile_m * TILE_M + mm) * N_DIM + (tile_n * TILE_N + tid)]; + __syncthreads(); + + float vals[TILE_M]; + float amax = 0.0f; + #pragma unroll + for (int mm = 0; mm < TILE_M; mm++) { + float v = (float)lds[mm * LDS_STRIDE + tid]; + vals[mm] = v; + amax = fmaxf(amax, fabsf(v)); + } + int e8 = (int)floorf(log2f(fmaxf(amax, 1e-38f))) + 127; + e8 = max(0, min(254, e8)); + float qscale = exp2f((float)(127 - e8)); + + const long long n = tile_n * TILE_N + tid; + __hip_fp8_storage_t out[TILE_M]; + #pragma unroll + for (int mm = 0; mm < TILE_M; mm++) + out[mm] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, vals[mm] * qscale)), __HIP_SATFINITE, __HIP_E4M3); + // 32 contiguous fp8 along M -> two 16-byte vector stores + long long obase = n * M_DIM + (long long)(tile_m * TILE_M); + *reinterpret_cast(&q[obase]) = *reinterpret_cast(&out[0]); + *reinterpret_cast(&q[obase + 16]) = *reinterpret_cast(&out[16]); + e8_out[n * (M_DIM / BLK) + tile_m] = (uint8_t)e8; +}