mirror of
https://github.com/Niko1221/Strata.git
synced 2026-11-09 03:59:45 -10:00
Merge commit 'b8d604d28' into f-clean-1a
This commit is contained in:
Niko1221 committed
commit
b0fc432ed9
126 files changed
+44033
-1557
No files matched your search
@@ -0,0 +1,218 @@
|
||||
name: Backend compile checks
|
||||
|
||||
# Compile and link every engine backend with a pinned toolchain, without a GPU: nothing is run, so a call into a kernel that only
|
||||
# exists for another backend (or an API a toolkit version dropped) fails here instead of on a user's machine. Hardware runs are separate.
|
||||
# CUDA 12.9 and 13.0 nvidia/cuda devel container targets strata, strata_kq (CMake preset cuda-compile)
|
||||
# HIP (ROCm 7.2) rocm/dev container target strata (CMake preset hip-compile)
|
||||
# SYCL (oneAPI 2026.1) ubuntu + Intel's apt repo target strata of sycl/ (the commands of packaging/build_sycl_linux.sh)
|
||||
# Metal macOS runner targets strata and strata_ds4 (metal/; metal-build.yml keeps its own strata build)
|
||||
# llama.cpp (ggml) comes at the commit pinned in setup.py (LLAMA_CPP_COMMIT), the same one the release builds use.
|
||||
# Architectures are fewer than the release package's (compile time); the flags that decide which sources and defines are built are the release's.
|
||||
|
||||
on:
|
||||
push:
|
||||
branches: [main]
|
||||
paths:
|
||||
- .github/workflows/backends.yml
|
||||
- CMakeLists.txt
|
||||
- CMakePresets.json
|
||||
- cmake/**
|
||||
- src/**
|
||||
- include/**
|
||||
- third_party/**
|
||||
- sycl/**
|
||||
- metal/**
|
||||
- setup.py
|
||||
- packaging/build_cuda_linux.sh
|
||||
- packaging/build_hip_linux.sh
|
||||
- packaging/build_sycl_linux.sh
|
||||
pull_request:
|
||||
paths:
|
||||
- .github/workflows/backends.yml
|
||||
- CMakeLists.txt
|
||||
- CMakePresets.json
|
||||
- cmake/**
|
||||
- src/**
|
||||
- include/**
|
||||
- third_party/**
|
||||
- sycl/**
|
||||
- metal/**
|
||||
- setup.py
|
||||
- packaging/build_cuda_linux.sh
|
||||
- packaging/build_hip_linux.sh
|
||||
- packaging/build_sycl_linux.sh
|
||||
workflow_dispatch:
|
||||
|
||||
permissions:
|
||||
contents: read
|
||||
|
||||
concurrency:
|
||||
group: backends-${{ github.ref }}
|
||||
cancel-in-progress: true
|
||||
|
||||
env:
|
||||
CCACHE_DIR: ${{ github.workspace }}/.ccache
|
||||
CCACHE_MAXSIZE: 2G
|
||||
|
||||
jobs:
|
||||
cuda:
|
||||
name: CUDA ${{ matrix.cuda }}
|
||||
runs-on: ubuntu-24.04
|
||||
timeout-minutes: 120
|
||||
container: nvidia/cuda:${{ matrix.image }}
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
include:
|
||||
# the cuda12 package is the one for Pascal and Volta: sm_60 needs STRATA_EXPERIMENTAL_SM60 (packaging/build_cuda_linux.sh)
|
||||
- { cuda: '12.9', image: '12.9.1-devel-ubuntu24.04', archs: '60-real;61-real;70-real;86-real', sm60: 'ON' }
|
||||
- { cuda: '13.0', image: '13.0.2-devel-ubuntu24.04', archs: '75-real;86-real;120-real', sm60: 'OFF' }
|
||||
steps:
|
||||
- name: Tools
|
||||
run: |
|
||||
apt-get update
|
||||
DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends git ca-certificates cmake ninja-build ccache
|
||||
cmake --version | head -1
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- name: llama.cpp (ggml) at the pinned commit
|
||||
run: |
|
||||
commit=$(sed -n 's/^LLAMA_CPP_COMMIT = "\(.*\)"/\1/p' setup.py)
|
||||
mkdir llama.cpp && cd llama.cpp && git init -q
|
||||
git remote add origin https://github.com/ggml-org/llama.cpp.git
|
||||
git fetch -q --depth 1 origin "$commit" && git checkout -q FETCH_HEAD
|
||||
- uses: actions/cache@v4
|
||||
with:
|
||||
path: .ccache
|
||||
key: ccache-cuda${{ matrix.cuda }}-${{ hashFiles('CMakeLists.txt', 'src/**', 'include/**', 'cmake/**') }}
|
||||
restore-keys: ccache-cuda${{ matrix.cuda }}-
|
||||
- name: Configure (preset cuda-compile)
|
||||
run: |
|
||||
nvcc --version | tail -2
|
||||
cmake --preset cuda-compile -B build -DSTRATA_GGML_DIR="$PWD/llama.cpp" \
|
||||
"-DCMAKE_CUDA_ARCHITECTURES=${{ matrix.archs }}" -DSTRATA_EXPERIMENTAL_SM60=${{ matrix.sm60 }} \
|
||||
-DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DCMAKE_CUDA_COMPILER_LAUNCHER=ccache
|
||||
- name: Build strata and strata_kq
|
||||
run: cmake --build build --target strata strata_kq -j "$(nproc)"
|
||||
- name: Cache statistics
|
||||
run: ccache -s | head -8
|
||||
|
||||
hip:
|
||||
name: HIP ROCm 7.2
|
||||
runs-on: ubuntu-24.04
|
||||
timeout-minutes: 120
|
||||
container: rocm/dev-ubuntu-24.04:7.2.4-complete
|
||||
steps:
|
||||
- name: Tools
|
||||
run: |
|
||||
apt-get update
|
||||
DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends git ca-certificates cmake ninja-build ccache
|
||||
cmake --version | head -1
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- name: llama.cpp (ggml) at the pinned commit
|
||||
run: |
|
||||
commit=$(sed -n 's/^LLAMA_CPP_COMMIT = "\(.*\)"/\1/p' setup.py)
|
||||
mkdir llama.cpp && cd llama.cpp && git init -q
|
||||
git remote add origin https://github.com/ggml-org/llama.cpp.git
|
||||
git fetch -q --depth 1 origin "$commit" && git checkout -q FETCH_HEAD
|
||||
- uses: actions/cache@v4
|
||||
with:
|
||||
path: .ccache
|
||||
key: ccache-hip72-${{ hashFiles('CMakeLists.txt', 'src/**', 'include/**', 'cmake/**') }}
|
||||
restore-keys: ccache-hip72-
|
||||
- name: Configure (preset hip-compile)
|
||||
run: |
|
||||
/opt/rocm/llvm/bin/clang++ --version | head -1
|
||||
cmake --preset hip-compile -B build -DSTRATA_GGML_DIR="$PWD/llama.cpp" \
|
||||
-DCMAKE_HIP_COMPILER=/opt/rocm/llvm/bin/clang++ -DCMAKE_HIP_COMPILER_ROCM_ROOT=/opt/rocm -DCMAKE_PREFIX_PATH=/opt/rocm \
|
||||
-DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DCMAKE_HIP_COMPILER_LAUNCHER=ccache
|
||||
- name: Build strata
|
||||
run: cmake --build build --target strata -j "$(nproc)"
|
||||
- name: Cache statistics
|
||||
run: ccache -s | head -8
|
||||
|
||||
sycl:
|
||||
# The engine needs oneAPI 2026.1 or newer (sycl/include/dpct/util.hpp uses the entangle and chunked_partition extensions, which the
|
||||
# 2025.3.2 compiler does not have: it stops at "no member named 'entangle'"; the intel/oneapi-basekit images end at 2025.3), so the
|
||||
# compiler comes from Intel's apt repository at a pinned version.
|
||||
name: SYCL oneAPI 2026.1
|
||||
runs-on: ubuntu-24.04
|
||||
timeout-minutes: 120
|
||||
container: ubuntu:24.04
|
||||
steps:
|
||||
- name: Tools and oneAPI 2026.1
|
||||
run: |
|
||||
apt-get update
|
||||
DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends git ca-certificates cmake ninja-build ccache wget gpg
|
||||
wget -qO- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor -o /usr/share/keyrings/oneapi-archive-keyring.gpg
|
||||
echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" > /etc/apt/sources.list.d/oneAPI.list
|
||||
apt-get update
|
||||
DEBIAN_FRONTEND=noninteractive apt-get install -y --no-install-recommends \
|
||||
intel-oneapi-compiler-dpcpp-cpp-2026.1=2026.1.1-325 intel-oneapi-mkl-devel-2026.1=2026.1.0-236
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- name: llama.cpp (ggml) at the pinned commit
|
||||
run: |
|
||||
commit=$(sed -n 's/^LLAMA_CPP_COMMIT = "\(.*\)"/\1/p' setup.py)
|
||||
mkdir llama.cpp && cd llama.cpp && git init -q
|
||||
git remote add origin https://github.com/ggml-org/llama.cpp.git
|
||||
git fetch -q --depth 1 origin "$commit" && git checkout -q FETCH_HEAD
|
||||
- uses: actions/cache@v4
|
||||
with:
|
||||
path: .ccache
|
||||
key: ccache-sycl-${{ hashFiles('CMakeLists.txt', 'src/**', 'include/**', 'sycl/**') }}
|
||||
restore-keys: ccache-sycl-
|
||||
- name: Configure and build (the commands of packaging/build_sycl_linux.sh)
|
||||
shell: bash
|
||||
run: |
|
||||
. /opt/intel/oneapi/setvars.sh --force >/dev/null
|
||||
icpx --version | head -1
|
||||
cmake -S sycl -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_COMPILER=icpx \
|
||||
-DSTRATA_SYCL_PARITY=OFF -DSTRATA_SYCL_IQ_DECODE_FAST=ON -DSTRATA_GGML_DIR="$PWD/llama.cpp" \
|
||||
-DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache
|
||||
cmake --build build --target strata -j "$(nproc)"
|
||||
- name: Cache statistics
|
||||
run: ccache -s | head -8
|
||||
|
||||
metal:
|
||||
name: Metal strata + strata_ds4
|
||||
# the macOS 26 SDK: metal/src/mtl_rt.mm names MTLGPUFamilyMetal4 behind @available (a run-time check); the names must be in the SDK
|
||||
runs-on: macos-26
|
||||
timeout-minutes: 60
|
||||
env:
|
||||
# metal/setup_mac.py's metal-cpp pin (the same as metal-build.yml)
|
||||
METAL_CPP_URL: https://developer.apple.com/metal/cpp/files/metal-cpp_26.zip
|
||||
METAL_CPP_SHA256: 4df3c078b9aadcb516212e9cb03004cbc5ce9a3e9c068fa3144d021db585a3a4
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- name: Tools
|
||||
run: brew install ninja ccache
|
||||
- name: metal-cpp
|
||||
run: |
|
||||
curl -fsSL "$METAL_CPP_URL" -o metal-cpp.zip
|
||||
echo "$METAL_CPP_SHA256 metal-cpp.zip" | shasum -a 256 -c -
|
||||
unzip -q metal-cpp.zip -d _mcpp
|
||||
- name: llama.cpp (ggml) at the pinned commit
|
||||
run: |
|
||||
commit=$(sed -n 's/^LLAMA_CPP_COMMIT = "\(.*\)"/\1/p' setup.py)
|
||||
mkdir llama.cpp && cd llama.cpp && git init -q
|
||||
git remote add origin https://github.com/ggml-org/llama.cpp.git
|
||||
git fetch -q --depth 1 origin "$commit" && git checkout -q FETCH_HEAD
|
||||
- uses: actions/cache@v4
|
||||
with:
|
||||
path: .ccache
|
||||
key: ccache-metal-${{ hashFiles('CMakeLists.txt', 'src/**', 'include/**', 'metal/**') }}
|
||||
restore-keys: ccache-metal-
|
||||
- name: Configure
|
||||
run: |
|
||||
cmake -S metal -B build-metal -G Ninja -DCMAKE_BUILD_TYPE=Release \
|
||||
-DMETAL_CPP_DIR="$PWD/_mcpp/metal-cpp" -DSTRATA_GGML_DIR="$PWD/llama.cpp" \
|
||||
-DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache
|
||||
- name: Build and link strata and strata_ds4
|
||||
run: cmake --build build-metal --target strata strata_ds4
|
||||
@@ -0,0 +1,185 @@
|
||||
name: Checks
|
||||
|
||||
# The fast checks (tools/check), one job per step so a failure names its step: the CPU CMake preset, the Python suites, the host's
|
||||
# typecheck / test / build. No GPU, no model and no download of weights. The config snapshot has its own workflow (snapshot.yml)
|
||||
# and the backend compile checks theirs (backends.yml). Run the same thing by hand: python tools/check --only <step>
|
||||
|
||||
on:
|
||||
push:
|
||||
branches: [main]
|
||||
paths:
|
||||
- .github/workflows/check.yml
|
||||
- tools/check
|
||||
- requirements*.txt
|
||||
- CMakeLists.txt
|
||||
- CMakePresets.json
|
||||
- cmake/**
|
||||
- src/**
|
||||
- include/**
|
||||
- bench/micro/**
|
||||
- third_party/**
|
||||
- tests/**
|
||||
- serve/**
|
||||
- strata.py
|
||||
- setup.py
|
||||
- strata_cli/**
|
||||
- strata_control/**
|
||||
- tools/**
|
||||
- packaging/**
|
||||
- docker/**
|
||||
- apps/host/**
|
||||
- apps/desktop/**
|
||||
- metal/**
|
||||
- sycl/**
|
||||
pull_request:
|
||||
paths:
|
||||
- .github/workflows/check.yml
|
||||
- tools/check
|
||||
- requirements*.txt
|
||||
- CMakeLists.txt
|
||||
- CMakePresets.json
|
||||
- cmake/**
|
||||
- src/**
|
||||
- include/**
|
||||
- bench/micro/**
|
||||
- third_party/**
|
||||
- tests/**
|
||||
- serve/**
|
||||
- strata.py
|
||||
- setup.py
|
||||
- strata_cli/**
|
||||
- strata_control/**
|
||||
- tools/**
|
||||
- packaging/**
|
||||
- docker/**
|
||||
- apps/host/**
|
||||
- apps/desktop/**
|
||||
- metal/**
|
||||
- sycl/**
|
||||
workflow_dispatch:
|
||||
|
||||
permissions:
|
||||
contents: read
|
||||
|
||||
concurrency:
|
||||
group: checks-${{ github.ref }}
|
||||
cancel-in-progress: true
|
||||
|
||||
jobs:
|
||||
linux:
|
||||
name: Linux / ${{ matrix.name }}
|
||||
runs-on: ubuntu-24.04
|
||||
timeout-minutes: 90
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
include:
|
||||
- { name: cpu CMake preset, only: cmake-cpu }
|
||||
- { name: Python tests/, only: py-tests }
|
||||
- { name: Python serve/, only: py-serve }
|
||||
- { name: Python tools/, only: py-tools }
|
||||
- { name: host typecheck + test + build, only: "host-typecheck,host-test,host-build" }
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- uses: actions/setup-python@v5
|
||||
with:
|
||||
python-version: '3.12'
|
||||
cache: pip
|
||||
cache-dependency-path: requirements*.txt
|
||||
- uses: actions/setup-node@v4
|
||||
if: ${{ startsWith(matrix.only, 'host') }}
|
||||
with:
|
||||
node-version: '24'
|
||||
cache: npm
|
||||
cache-dependency-path: apps/host/package-lock.json
|
||||
- name: Python packages
|
||||
run: python -m pip install -r requirements.txt -r requirements-dev.txt -r requirements-build.txt
|
||||
- name: tools/check
|
||||
run: python tools/check --only "${{ matrix.only }}" --keep --scratch "$RUNNER_TEMP/check" --json "$RUNNER_TEMP/check/result.json"
|
||||
- name: Logs
|
||||
if: ${{ failure() }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: check-linux-${{ strategy.job-index }}
|
||||
path: |
|
||||
${{ runner.temp }}/check/*.log
|
||||
${{ runner.temp }}/check/result.json
|
||||
retention-days: 7
|
||||
|
||||
windows:
|
||||
name: Windows / ${{ matrix.name }}
|
||||
runs-on: windows-2022
|
||||
timeout-minutes: 120
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
include:
|
||||
- { name: cpu CMake preset (MSVC), only: cmake-cpu }
|
||||
- { name: Python tests/, only: py-tests }
|
||||
- { name: Python serve/, only: py-serve }
|
||||
- { name: Python tools/, only: py-tools }
|
||||
- { name: host typecheck + test + build, only: "host-typecheck,host-test,host-build" }
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- uses: actions/setup-python@v5
|
||||
with:
|
||||
python-version: '3.12'
|
||||
cache: pip
|
||||
cache-dependency-path: requirements*.txt
|
||||
- uses: actions/setup-node@v4
|
||||
if: ${{ startsWith(matrix.only, 'host') }}
|
||||
with:
|
||||
node-version: '24'
|
||||
cache: npm
|
||||
cache-dependency-path: apps/host/package-lock.json
|
||||
- name: Python packages
|
||||
run: python -m pip install -r requirements.txt -r requirements-dev.txt -r requirements-build.txt
|
||||
# tools/check finds Visual Studio's compiler itself (vcvars64.bat); bash is Git for Windows' (the tester-README tests need it)
|
||||
- name: tools/check
|
||||
run: python tools/check --only "${{ matrix.only }}" --keep --scratch "$env:RUNNER_TEMP\check" --json "$env:RUNNER_TEMP\check\result.json"
|
||||
- name: Logs
|
||||
if: ${{ failure() }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: check-windows-${{ strategy.job-index }}
|
||||
path: |
|
||||
${{ runner.temp }}\check\*.log
|
||||
${{ runner.temp }}\check\result.json
|
||||
retention-days: 7
|
||||
|
||||
macos:
|
||||
# Cheap and useful (the audit found the host and Python fixtures reading the Mac's own paths), but this job has not run on a
|
||||
# Mac runner yet: it reports and does not block until its first green run.
|
||||
name: macOS / ${{ matrix.name }}
|
||||
runs-on: macos-15
|
||||
timeout-minutes: 60
|
||||
continue-on-error: true
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
include:
|
||||
- { name: Python tools/, only: py-tools }
|
||||
- { name: host typecheck + test + build, only: "host-typecheck,host-test,host-build" }
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- uses: actions/setup-python@v5
|
||||
with:
|
||||
python-version: '3.12'
|
||||
cache: pip
|
||||
cache-dependency-path: requirements*.txt
|
||||
- uses: actions/setup-node@v4
|
||||
if: ${{ startsWith(matrix.only, 'host') }}
|
||||
with:
|
||||
node-version: '24'
|
||||
cache: npm
|
||||
cache-dependency-path: apps/host/package-lock.json
|
||||
- name: Python packages
|
||||
run: python -m pip install -r requirements.txt -r requirements-dev.txt
|
||||
- name: tools/check
|
||||
run: python tools/check --only "${{ matrix.only }}"
|
||||
@@ -0,0 +1,83 @@
|
||||
name: Config snapshot
|
||||
|
||||
# tests/snapshot: what `strata launch`, the app's plain Load and the app's Models-screen Load hand the engine (arguments, environment, drafter,
|
||||
# confidence), for every simulated machine class and model set, against the golden files in tests/snapshot/golden. Nothing is loaded,
|
||||
# downloaded or run on a GPU. Windows and Linux must give byte-identical output, so both run. A change that moves what users
|
||||
# get has to update the golden files in the same commit (STRATA_SNAPSHOT_UPDATE=1 python -m pytest tests/snapshot, see tests/snapshot/README.md).
|
||||
|
||||
on:
|
||||
push:
|
||||
branches: [main]
|
||||
paths:
|
||||
- .github/workflows/snapshot.yml
|
||||
- tests/snapshot/**
|
||||
- serve/**
|
||||
- strata_cli/**
|
||||
- strata_control/**
|
||||
- setup.py
|
||||
- strata.py
|
||||
- requirements*.txt
|
||||
- apps/host/server/**
|
||||
- apps/host/package*.json
|
||||
- packaging/**
|
||||
- tools/check
|
||||
pull_request:
|
||||
paths:
|
||||
- .github/workflows/snapshot.yml
|
||||
- tests/snapshot/**
|
||||
- serve/**
|
||||
- strata_cli/**
|
||||
- strata_control/**
|
||||
- setup.py
|
||||
- strata.py
|
||||
- requirements*.txt
|
||||
- apps/host/server/**
|
||||
- apps/host/package*.json
|
||||
- packaging/**
|
||||
- tools/check
|
||||
workflow_dispatch:
|
||||
|
||||
permissions:
|
||||
contents: read
|
||||
|
||||
concurrency:
|
||||
group: snapshot-${{ github.ref }}
|
||||
cancel-in-progress: true
|
||||
|
||||
jobs:
|
||||
snapshot:
|
||||
name: ${{ matrix.os }}
|
||||
runs-on: ${{ matrix.os }}
|
||||
timeout-minutes: 30
|
||||
strategy:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
os: [ubuntu-24.04, windows-2022]
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
with:
|
||||
persist-credentials: false
|
||||
- uses: actions/setup-python@v5
|
||||
with:
|
||||
python-version: '3.12'
|
||||
cache: pip
|
||||
cache-dependency-path: requirements*.txt
|
||||
- uses: actions/setup-node@v4
|
||||
with:
|
||||
node-version: '24'
|
||||
cache: npm
|
||||
cache-dependency-path: apps/host/package-lock.json
|
||||
- name: Python packages
|
||||
run: python -m pip install -r requirements.txt -r requirements-dev.txt
|
||||
- name: Host packages (the app's Load runs in Node)
|
||||
working-directory: apps/host
|
||||
run: npm ci --no-audit --no-fund
|
||||
- name: tools/check snapshot
|
||||
run: python tools/check --only snapshot --keep --scratch "${{ runner.temp }}/check"
|
||||
- name: Logs
|
||||
if: ${{ failure() }}
|
||||
uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: snapshot-${{ matrix.os }}
|
||||
path: ${{ runner.temp }}/check/*.log
|
||||
retention-days: 7
|
||||
+71
-18
@@ -1,12 +1,14 @@
|
||||
# Strata Engine - top-level build (P1.S1).
|
||||
# Strata Engine - top-level build.
|
||||
#
|
||||
# Targets (phase-1 §P1.S1): `strata_artifact` (no CUDA dependency), `tests`. `strata_plan` and the C++
|
||||
# `strata-pack` are added when their sources exist - a target with no sources would make the build graph
|
||||
# claim something the tree does not have, which is worse than an absent target.
|
||||
# One project builds the engine for every backend: CUDA (STRATA_ENABLE_CUDA), HIP (STRATA_ENABLE_HIP; the layer-split
|
||||
# and DeepSeek engines are the STRATA_BUILD_SPLIT / STRATA_BUILD_DSV41 options), the CPU-only build, and the test
|
||||
# programs. The Intel Arc (SYCL) engine is a separate project in sycl/ (STRATA_ENABLE_SYCL), and the Metal engine
|
||||
# one in metal/. CMakePresets.json has the CPU check preset (`cmake --preset cpu-check`) and the compile-only
|
||||
# backend presets used by tools/check and the CI workflows; the release scripts in packaging/ pass their own flags.
|
||||
#
|
||||
# CUDA targets arrive in Phase 2. The kernels need sm_80 or newer (tf32 mma, bf16 math) and the guard
|
||||
# below refuses anything older: silently building for a pre-Ampere arch would produce a binary that runs
|
||||
# and is wrong about its own performance, which is the worst of both outcomes.
|
||||
# The CUDA kernels need sm_80 or newer (tf32 mma, bf16 math) and the guard below refuses anything older: silently
|
||||
# building for a pre-Ampere arch would produce a binary that runs and is wrong about its own performance, which is
|
||||
# the worst of both outcomes.
|
||||
cmake_minimum_required(VERSION 3.24)
|
||||
project(strata VERSION 0.1.58 LANGUAGES CXX)
|
||||
# EXPERIMENTAL: the Intel Arc engine, the SYCL port in sycl/ (PR #423; icpx + oneMKL, docs/INTEL_ARC.md). It is its
|
||||
@@ -158,17 +160,28 @@ endif()
|
||||
if(STRATA_ENABLE_HIP)
|
||||
include(cmake/hip_backend.cmake)
|
||||
endif()
|
||||
# The test suite and the micro benchmarks live in tests/ and bench/micro/, which the published source leaves out.
|
||||
# Tests are on by default when the build can run them. What they need: their sources (src/**/*_test.cpp, tests/core,
|
||||
# bench/micro; a source drop without them has nothing to build) and the x86 CPU kernels, which most of the CPU suite drives
|
||||
# (arm64 stays off by default: -DSTRATA_BUILD_TESTS=ON still builds what builds there). GPU tests are registered only
|
||||
# when STRATA_ENABLE_CUDA / STRATA_ENABLE_HIP is on, and the configure summary below says which kinds were registered, so
|
||||
# a missing GPU is a stated skip and never a silent pass. Release builds say -DSTRATA_BUILD_TESTS=OFF themselves
|
||||
# (packaging/, setup.py, Dockerfile, spark/).
|
||||
set(_strata_tests_default OFF)
|
||||
if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/tests/CMakeLists.txt AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/bench/micro)
|
||||
if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/tests/core AND EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/bench/micro AND STRATA_CPU_X86)
|
||||
set(_strata_tests_default ON)
|
||||
endif()
|
||||
if(NOT STRATA_CPU_X86)
|
||||
# arm64 Linux (DGX Spark): much of the suite drives the x86 CPU kernels, which this build does not have. A default, not a
|
||||
# prohibition: -DSTRATA_BUILD_TESTS=ON still builds what builds (classifying the tests by capability is a follow-up).
|
||||
set(_strata_tests_default OFF)
|
||||
endif()
|
||||
option(STRATA_BUILD_TESTS "Build the test suite" ${_strata_tests_default})
|
||||
# A test that needs a model file (a real expert blob from a pack, which is not in the source tree): registered as the real test
|
||||
# when the file is there, and otherwise as a stated skip that ctest lists as "Skipped" with this reason, never as a pass of nothing
|
||||
# and not as a failure on a machine that has no pack.
|
||||
function(strata_add_pack_test name needs)
|
||||
if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${needs}")
|
||||
add_test(NAME ${name} COMMAND ${ARGN} WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
else()
|
||||
add_test(NAME ${name} COMMAND ${CMAKE_COMMAND} -E echo "SKIPPED: ${name} needs ${needs} (a Q2_0 expert pack, a model file), which is not in this tree")
|
||||
set_tests_properties(${name} PROPERTIES SKIP_REGULAR_EXPRESSION "SKIPPED: ")
|
||||
endif()
|
||||
endfunction()
|
||||
option(STRATA_BUILD_CONVERSATION_TESTS "Build standalone conversation cache tests" OFF)
|
||||
option(STRATA_WERROR "Treat compiler warnings as errors" OFF)
|
||||
# #2027: defined here, before the GPU section reads it to register the native-expert tests (it was defined at the
|
||||
@@ -1045,6 +1058,18 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
|
||||
add_executable(ds4_route_guard_test EXCLUDE_FROM_ALL tests/dsv41/ds4_route_guard_test.cu src/dsv41/ds4_kernels.cu)
|
||||
target_include_directories(ds4_route_guard_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include ${_strata_gpu_include_directories})
|
||||
target_link_libraries(ds4_route_guard_test PRIVATE strata_dsv41_kernels ${_strata_gpu_runtime_target} Threads::Threads)
|
||||
# the DeepSeek sampler's statistics harness as a CUDA standalone (cleanup 4.4): the same src/dsv41/sampler_stat.inc that `strata-dsv41
|
||||
# --sampler-test` and the Metal engine compile, here with the CUDA sampler kernel, so the two backends' picks can be compared sample by
|
||||
# sample (same Philox counters). It includes vendor/sampler.cu itself, so it links no engine library. NVIDIA only; needs a GPU, no
|
||||
# model; built by name: cmake --build <dir> --target ds4_sampler_stat
|
||||
# usage: ds4_sampler_stat ROWS.bin N_VOCAB N_SAMPLES OUTPREFIX (see the header of src/dsv41/sampler_stat.inc for the files it writes)
|
||||
if(NOT STRATA_ENABLE_HIP)
|
||||
add_executable(ds4_sampler_stat EXCLUDE_FROM_ALL src/dsv41/sampler_stat_main.cu)
|
||||
target_include_directories(ds4_sampler_stat PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include ${CMAKE_CURRENT_SOURCE_DIR}/src/dsv41
|
||||
${_strata_gpu_include_directories})
|
||||
target_compile_options(ds4_sampler_stat PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:SHELL:-include ds4_launch.h>)
|
||||
target_link_libraries(ds4_sampler_stat PRIVATE ${_strata_gpu_runtime_target})
|
||||
endif()
|
||||
if(STRATA_NATIVE_EXPERTS)
|
||||
# the trellis experts (device kernels ds4_tq.cu, CPU lane ds4_tq_cpu.cpp) against a double reference; needs a GPU, no model
|
||||
# both test .cpp files launch kernels, so they compile as HIP on the AMD backend and as CUDA on the NVIDIA one
|
||||
@@ -2064,6 +2089,17 @@ else()
|
||||
src/kernels/cpu/expert_layout.cpp src/kernels/cpu/hadamard.cpp)
|
||||
endif()
|
||||
target_include_directories(strata_kernels_cpu PUBLIC include)
|
||||
# expert_layout.cpp (above) names the AVX1 rung of the Q2_0 rows and the activation quantizer; with STRATA_NATIVE_EXPERTS=ON
|
||||
# q2_avx1.cpp joins the library in the native-experts section below (default, release builds). Without it the library has
|
||||
# to carry that file itself, or every program that links it fails with undefined *_avx1 symbols.
|
||||
if(STRATA_CPU_X86 AND NOT STRATA_NATIVE_EXPERTS)
|
||||
target_sources(strata_kernels_cpu PRIVATE src/kernels/cpu/q2_avx1.cpp)
|
||||
if(MSVC)
|
||||
set_source_files_properties(src/kernels/cpu/q2_avx1.cpp PROPERTIES COMPILE_OPTIONS "/arch:AVX")
|
||||
else()
|
||||
set_source_files_properties(src/kernels/cpu/q2_avx1.cpp PROPERTIES COMPILE_OPTIONS "-msse4.2;-mavx;-mno-avx2;-mno-fma;-mno-f16c")
|
||||
endif()
|
||||
endif()
|
||||
# Plan v0.3 P6: only the kernel files get the AVX-512 / AVX2 flags; the pool and the probe run on any x86-64,
|
||||
# and the engine picks the kernels at run time (cpu_avx512_ok).
|
||||
if(MSVC)
|
||||
@@ -2894,12 +2930,11 @@ if(STRATA_BUILD_TESTS)
|
||||
target_link_libraries(expert_parity PRIVATE strata_kernels_cpu)
|
||||
# The pack is a REQUIRED input, not an optional one: this test compares against a real expert blob, and a
|
||||
# missing file would otherwise be reported as a pass by a test that skipped everything.
|
||||
add_test(NAME expert_parity COMMAND expert_parity --selftest
|
||||
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
strata_add_pack_test(expert_parity pack/full/experts.bin expert_parity --selftest)
|
||||
|
||||
add_executable(pool_test src/kernels/cpu/pool_test.cpp)
|
||||
target_link_libraries(pool_test PRIVATE strata_kernels_cpu)
|
||||
add_test(NAME pool_test COMMAND pool_test --selftest WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR})
|
||||
strata_add_pack_test(pool_test pack/full/experts.bin pool_test --selftest)
|
||||
# where the workers are placed (one per physical core by default) and --pool-cores (real pools, no model)
|
||||
add_executable(pool_cores_test src/kernels/cpu/pool_cores_test.cpp)
|
||||
target_link_libraries(pool_cores_test PRIVATE strata_kernels_cpu)
|
||||
@@ -3059,8 +3094,26 @@ if(STRATA_BUILD_DSV41 AND NOT STRATA_ENABLE_CUDA AND NOT STRATA_ENABLE_HIP AND N
|
||||
endif()
|
||||
|
||||
# ---- strata-split: the opt-in engine that splits every layer across 2-4 NVIDIA cards (src/split, from sFNX, MIT, John Fink) ----
|
||||
# OFF by default, Linux + CUDA only, no other target depends on it; see src/split/CMakeLists.txt and docs/SPLIT_ENGINE.md.
|
||||
# OFF by default, Linux + CUDA or HIP only, no other target depends on it; see src/split/CMakeLists.txt and docs/SPLIT_ENGINE.md.
|
||||
option(STRATA_BUILD_SPLIT "Build strata-split: sFNX's split-every-layer engine, host-memory exchange (Linux, CUDA or HIP gfx1201, 2-4 cards)" OFF)
|
||||
if(STRATA_BUILD_SPLIT)
|
||||
add_subdirectory(src/split)
|
||||
endif()
|
||||
|
||||
# What this configure registered, so a missing capability is a stated skip (tools/check prints this line too).
|
||||
if(STRATA_BUILD_TESTS)
|
||||
set(_strata_test_kinds "cpu")
|
||||
if(STRATA_ENABLE_CUDA)
|
||||
string(APPEND _strata_test_kinds ", cuda")
|
||||
elseif(STRATA_ENABLE_HIP)
|
||||
string(APPEND _strata_test_kinds ", hip")
|
||||
else()
|
||||
string(APPEND _strata_test_kinds " (gpu tests not registered: neither STRATA_ENABLE_CUDA nor STRATA_ENABLE_HIP is on)")
|
||||
endif()
|
||||
if(NOT STRATA_NATIVE_EXPERTS)
|
||||
string(APPEND _strata_test_kinds " (native-expert tests not registered: STRATA_NATIVE_EXPERTS=OFF)")
|
||||
endif()
|
||||
message(STATUS "strata tests: ON - ${_strata_test_kinds}")
|
||||
else()
|
||||
message(STATUS "strata tests: OFF (STRATA_BUILD_TESTS=OFF)")
|
||||
endif()
|
||||
@@ -0,0 +1,78 @@
|
||||
{
|
||||
"version": 5,
|
||||
"cmakeMinimumRequired": { "major": 3, "minor": 24, "patch": 0 },
|
||||
"configurePresets": [
|
||||
{
|
||||
"name": "cpu-check",
|
||||
"displayName": "CPU check build (no GPU, no network)",
|
||||
"description": "The fast check: the CPU-only engine pieces and every test that needs no GPU and no model. STRATA_NATIVE_EXPERTS=OFF, so ggml is not needed. tools/check runs this preset.",
|
||||
"generator": "Ninja",
|
||||
"binaryDir": "${sourceDir}/build/${presetName}",
|
||||
"cacheVariables": {
|
||||
"CMAKE_BUILD_TYPE": "Release",
|
||||
"STRATA_BUILD_TESTS": "ON",
|
||||
"STRATA_NATIVE_EXPERTS": "OFF"
|
||||
}
|
||||
},
|
||||
{
|
||||
"name": "cpu-check-native",
|
||||
"displayName": "CPU check build with the native expert tests",
|
||||
"description": "As cpu-check plus the native (IQ) expert kernels and their tests. Needs ggml: pass -DSTRATA_GGML_DIR=<llama.cpp checkout>, or let CMake fetch it.",
|
||||
"inherits": "cpu-check",
|
||||
"cacheVariables": { "STRATA_NATIVE_EXPERTS": "ON" }
|
||||
},
|
||||
{
|
||||
"name": "cuda-compile",
|
||||
"displayName": "CUDA compile-only (the release flags of packaging/build_cuda_linux.sh, no tests)",
|
||||
"description": "Builds the strata and strata_kq targets for the given CUDA architectures. No GPU is needed to compile; nothing is run. Pass -DSTRATA_GGML_DIR=<llama.cpp checkout> (CI does).",
|
||||
"generator": "Ninja",
|
||||
"binaryDir": "${sourceDir}/build/${presetName}",
|
||||
"cacheVariables": {
|
||||
"CMAKE_BUILD_TYPE": "Release",
|
||||
"STRATA_ENABLE_CUDA": "ON",
|
||||
"STRATA_BUILD_TESTS": "OFF",
|
||||
"STRATA_NATIVE_EXPERTS": "ON",
|
||||
"STRATA_PORTABLE": "ON",
|
||||
"STRATA_Q6K_EXPERTS": "ON",
|
||||
"STRATA_MMQ_KQUANTS": "ON",
|
||||
"STRATA_EXPERIMENTAL_SM60": "OFF",
|
||||
"CMAKE_CUDA_ARCHITECTURES": "86-real;89-real;89-virtual"
|
||||
}
|
||||
},
|
||||
{
|
||||
"name": "hip-compile",
|
||||
"displayName": "HIP compile-only (the release flags of packaging/build_hip_linux.sh, no tests)",
|
||||
"description": "Builds the strata target for gfx1100 and gfx1201 with ROCm's hipcc. Pass -DSTRATA_GGML_DIR=<llama.cpp checkout> and, if hipcc is not on PATH, -DCMAKE_HIP_COMPILER=<rocm>/llvm/bin/clang++.",
|
||||
"generator": "Ninja",
|
||||
"binaryDir": "${sourceDir}/build/${presetName}",
|
||||
"cacheVariables": {
|
||||
"CMAKE_BUILD_TYPE": "Release",
|
||||
"STRATA_ENABLE_HIP": "ON",
|
||||
"STRATA_ENABLE_CUDA": "OFF",
|
||||
"STRATA_BUILD_TESTS": "OFF",
|
||||
"STRATA_PREFILL_MMQ": "ON",
|
||||
"STRATA_PORTABLE": "ON",
|
||||
"CMAKE_HIP_ARCHITECTURES": "gfx1100;gfx1201"
|
||||
}
|
||||
}
|
||||
],
|
||||
"buildPresets": [
|
||||
{ "name": "cpu-check", "configurePreset": "cpu-check", "jobs": 0 },
|
||||
{ "name": "cpu-check-native", "configurePreset": "cpu-check-native", "jobs": 0 },
|
||||
{ "name": "cuda-compile", "configurePreset": "cuda-compile", "targets": ["strata", "strata_kq"], "jobs": 0 },
|
||||
{ "name": "hip-compile", "configurePreset": "hip-compile", "targets": ["strata"], "jobs": 0 }
|
||||
],
|
||||
"testPresets": [
|
||||
{
|
||||
"name": "cpu-check",
|
||||
"configurePreset": "cpu-check",
|
||||
"output": { "outputOnFailure": true, "shortProgress": true },
|
||||
"execution": { "timeout": 120, "noTestsAction": "error", "jobs": 4 }
|
||||
},
|
||||
{
|
||||
"name": "cpu-check-native",
|
||||
"configurePreset": "cpu-check-native",
|
||||
"inherits": "cpu-check"
|
||||
}
|
||||
]
|
||||
}
|
||||
@@ -11,6 +11,8 @@ import {tmpdir} from 'node:os';
|
||||
import {join} from 'node:path';
|
||||
const scratch=mkdtempSync(join(tmpdir(),'strata-host-suite-profile-'));
|
||||
const env={...process.env,STRATA_SETTINGS_FILE:join(scratch,'app-settings.json'),STRATA_COMPONENTS_DIR:join(scratch,'components'),LOCALAPPDATA:join(scratch,'Local'),APPDATA:join(scratch,'Roaming'),XDG_DATA_HOME:join(scratch,'xdg-data'),XDG_STATE_HOME:join(scratch,'xdg-state'),XDG_CONFIG_HOME:join(scratch,'xdg-config')};
|
||||
const child=spawn(process.execPath,['--test',...files],{cwd:root,stdio:'inherit',env});
|
||||
/* A test that never finishes fails after STRATA_TEST_TIMEOUT_MS (default 10 minutes) instead of holding the whole run for ever; extra arguments (a reporter, --test-name-pattern) go to node --test. */
|
||||
const timeout=Number(process.env.STRATA_TEST_TIMEOUT_MS)||600000;
|
||||
const child=spawn(process.execPath,['--test',`--test-timeout=${timeout}`,...process.argv.slice(2),...files],{cwd:root,stdio:'inherit',env});
|
||||
child.on('error',error=>{console.error(error.message);process.exitCode=1;});
|
||||
child.on('exit',(code,signal)=>{process.exitCode=code??(signal?1:0);});
|
||||
@@ -1,8 +1,11 @@
|
||||
// The chat's own folder and opening its files by path.
|
||||
import test from 'node:test';
|
||||
import assert from 'node:assert/strict';
|
||||
import { mkdtempSync, mkdirSync, rmSync, writeFileSync, symlinkSync, existsSync, readFileSync } from 'node:fs';
|
||||
import { tmpdir } from 'node:os';
|
||||
import { mkdtempSync, mkdirSync, rmSync, writeFileSync, symlinkSync, existsSync, readFileSync, realpathSync } from 'node:fs';
|
||||
import { tmpdir as osTmpdir } from 'node:os';
|
||||
// the canonical scratch root: on a Mac the temp folder is /var/folders/..., which the folder policy refuses as a system folder and which
|
||||
// is /private/var/... once resolved; the tests make their folders under the resolved name (the policy itself is unchanged)
|
||||
const tmpdir = () => realpathSync(osTmpdir());
|
||||
import { join } from 'node:path';
|
||||
import { createChatFolders, openFile, openCommand, resolveChatFile, refusedFolder } from '../server/chat-files.mjs';
|
||||
import { startScriptedWorker } from './helpers/scripted-worker.mjs';
|
||||
|
||||
@@ -68,7 +68,7 @@ test('STRATA_CLOUDFLARED picks exactly one file; STRATA_CLOUDFLARED_SEARCH=own k
|
||||
assert.deepEqual(own.map(c => c.source), ['strata', 'component']);
|
||||
});
|
||||
|
||||
test('a real executable on PATH is found and run (a user-local copy, never a system install)', { skip: process.platform === 'win32' }, async () => {
|
||||
test('a real executable on PATH is found and run (a user-local copy, never a system install)', { skip: process.platform === 'win32' && 'needs a POSIX executable file to run as cloudflared (not on Windows)' }, async () => {
|
||||
const dir = await mkdtemp(join(tmpdir(), 'strata-cf-path-'));
|
||||
try {
|
||||
const bin = join(dir, 'bin'); await mkdir(bin);
|
||||
|
||||
@@ -313,7 +313,8 @@ test('packaged: an app update that ends the app is recorded as finished by the n
|
||||
});
|
||||
|
||||
test('packaged: an update job an OLDER host left running is finished when the app program file is newer than the job', async () => {
|
||||
const exe = new URL('../../strata-desktop.exe', import.meta.url); // <install>/strata-desktop.exe as seen from <install>/host
|
||||
// <install>/strata-desktop(.exe) as seen from <install>/host; native-control.mjs looks for the name of the platform it runs on
|
||||
const exe = new URL(process.platform === 'win32' ? '../../strata-desktop.exe' : '../../strata-desktop', import.meta.url);
|
||||
const f = fixture({ withApp: true });
|
||||
try {
|
||||
const { fileURLToPath } = await import('node:url');
|
||||
|
||||
@@ -22,10 +22,22 @@ test('corrupt or unverified downloads cannot become executable', () => {
|
||||
assert.throws(()=>verifyConnector(bytes,null),/could not be verified/);
|
||||
});
|
||||
|
||||
test('failed concurrent setup shares one download, cleans up and can be retried', async () => {
|
||||
/** Waits until the fake download service has been asked; fails (instead of spinning for ever) when it never is. */
|
||||
async function untilAsked(isAsked){
|
||||
const deadline=Date.now()+10000;
|
||||
while(!isAsked()){
|
||||
if(Date.now()>deadline)throw new Error('the installer never asked the download service (it found a connector, or stopped before the request)');
|
||||
await new Promise(resolve=>setTimeout(resolve,5));
|
||||
}
|
||||
}
|
||||
|
||||
// The installer first looks for a connector that is already on the computer. The look is pinned to Strata's own empty folder
|
||||
// (search:false, an empty environment): a cloudflared installed on the machine that runs the test (PATH, /usr/local/bin, ~/.local/bin)
|
||||
// would otherwise be used, nothing would be downloaded, and the test would wait for ever for a request that never comes.
|
||||
test('failed concurrent setup shares one download, cleans up and can be retried', {timeout:30000}, async () => {
|
||||
const directory=await mkdtemp(join(tmpdir(),'strata-connector-test-'));
|
||||
let requests=0, release;
|
||||
const installer=createConnectorInstaller({directory,platform:'linux',arch:'x64',fetchImpl:async()=>{
|
||||
const installer=createConnectorInstaller({directory,platform:'linux',arch:'x64',env:{},home:directory,search:false,fetchImpl:async()=>{
|
||||
requests++;
|
||||
await new Promise(resolve=>{release=resolve;});
|
||||
return {ok:false};
|
||||
@@ -33,7 +45,7 @@ test('failed concurrent setup shares one download, cleans up and can be retried'
|
||||
try {
|
||||
const first=installer.ensure(),second=installer.ensure();
|
||||
assert.equal(first,second);
|
||||
while(!release)await new Promise(resolve=>setImmediate(resolve));
|
||||
await untilAsked(()=>release);
|
||||
release();
|
||||
const results=await Promise.allSettled([first,second]);
|
||||
assert.ok(results.every(result=>result.status==='rejected'));
|
||||
@@ -41,7 +53,7 @@ test('failed concurrent setup shares one download, cleans up and can be retried'
|
||||
assert.deepEqual(await readdir(directory),[]);
|
||||
release=null;
|
||||
const retry=installer.ensure();
|
||||
while(!release)await new Promise(resolve=>setImmediate(resolve));
|
||||
await untilAsked(()=>release);
|
||||
release();
|
||||
await assert.rejects(retry,/download service/);
|
||||
assert.equal(requests,2);
|
||||
|
||||
@@ -42,7 +42,7 @@ test('Windows uses a PowerShell FolderBrowserDialog passed as an encoded command
|
||||
assert.equal(nodialog.unavailable, true); assert.match(nodialog.reason, /Type the folder path/);
|
||||
});
|
||||
|
||||
test('the generated PowerShell parses without errors', { skip: process.platform !== 'win32' }, () => {
|
||||
test('the generated PowerShell parses without errors', { skip: process.platform !== 'win32' && 'needs Windows PowerShell to parse the script (Windows only)' }, () => {
|
||||
const shell = join(process.env.SystemRoot, 'System32', 'WindowsPowerShell', 'v1.0', 'powershell.exe');
|
||||
const out = execFileSync(shell,
|
||||
['-NoProfile', '-NonInteractive', '-Command', '$e=$null; [void][System.Management.Automation.Language.Parser]::ParseInput($env:SCRIPT,[ref]$null,[ref]$e); $e.Count'],
|
||||
|
||||
@@ -116,7 +116,7 @@ test('staging refuses a module hardlinked to another copy', t => {
|
||||
});
|
||||
|
||||
test('patching a disposable copy leaves the installed original hash unchanged', {
|
||||
skip: !process.env.STRATA_TEST_PI_AI_DIR,
|
||||
skip: !process.env.STRATA_TEST_PI_AI_DIR && 'STRATA_TEST_PI_AI_DIR is not set (the Harness AI package folder)',
|
||||
}, t => {
|
||||
const installed = join(process.env.STRATA_TEST_PI_AI_DIR, 'dist', 'api', 'simple-options.js');
|
||||
const before = sha256(readFileSync(installed));
|
||||
|
||||
@@ -43,6 +43,6 @@ test('missingHarnessImports names a package a plugin imports that is not staged'
|
||||
} finally { rmSync(root, { recursive: true, force: true }); }
|
||||
});
|
||||
|
||||
test('a staged Harness (STRATA_HARNESS_DIR) has every package its plugins import', { skip: !process.env.STRATA_HARNESS_DIR }, () => {
|
||||
test('a staged Harness (STRATA_HARNESS_DIR) has every package its plugins import', { skip: !process.env.STRATA_HARNESS_DIR && 'STRATA_HARNESS_DIR is not set to a built Harness' }, () => {
|
||||
assert.deepEqual(missingHarnessImports(process.env.STRATA_HARNESS_DIR), []);
|
||||
});
|
||||
@@ -12,11 +12,13 @@ import { join, resolve } from 'node:path';
|
||||
import { build } from 'esbuild';
|
||||
|
||||
const root = resolve(import.meta.dirname, '..');
|
||||
// the text arrives as JSON in argv[2], so no shell or platform touches its line ends
|
||||
// the text arrives as JSON in a file named by argv[2], so no shell or platform touches its line ends (and a long text is not cut by
|
||||
// the command-line length limit: Windows allows about 32 KB, which the 1500-section case below passes)
|
||||
const entry = [
|
||||
"import React from 'react'; import { renderToStaticMarkup } from 'react-dom/server';",
|
||||
"import { Markdown } from './src/components/Markdown.tsx'; import { splitPreview } from './src/lib/preview-render.ts';",
|
||||
'const text = JSON.parse(process.argv[2]);',
|
||||
"import { readFileSync } from 'node:fs';",
|
||||
"const text = JSON.parse(readFileSync(process.argv[2], 'utf8'));",
|
||||
'console.log(JSON.stringify({ html: renderToStaticMarkup(React.createElement(Markdown, { text })), slices: splitPreview(text).map(s => s.text) }));',
|
||||
].join('\n');
|
||||
const bundle = await build({
|
||||
@@ -30,8 +32,11 @@ const bundle = await build({
|
||||
const dir = mkdtempSync(join(tmpdir(), 'strata-crlf-'));
|
||||
const file = join(dir, 'render.cjs');
|
||||
writeFileSync(file, bundle.outputFiles[0].text);
|
||||
let renders = 0;
|
||||
function render(text) {
|
||||
const run = spawnSync(process.execPath, ['--max-old-space-size=200', file, JSON.stringify(text)], { cwd: root, encoding: 'utf8', timeout: 20000, maxBuffer: 20 * 1024 * 1024 });
|
||||
const input = join(dir, `text-${renders++}.json`);
|
||||
writeFileSync(input, JSON.stringify(text));
|
||||
const run = spawnSync(process.execPath, ['--max-old-space-size=200', file, input], { cwd: root, encoding: 'utf8', timeout: 20000, maxBuffer: 20 * 1024 * 1024 });
|
||||
return { ok: run.status === 0 && !run.error, out: run.stdout, err: (run.stderr || '') + (run.error?.message ?? ''), status: run.status, signal: run.signal };
|
||||
}
|
||||
|
||||
|
||||
@@ -66,7 +66,7 @@ test('choosing a folder that is gone or cannot be written says so in words', asy
|
||||
assert.throws(() => f.control.setDownloadDirectory(join(f.repo, 'setup.py')), /Choose a folder, not a file\./);
|
||||
});
|
||||
|
||||
test('Windows: a drive picked as D:\\ becomes D:\\Strata, and the person is not asked for a different folder', { skip: process.platform !== 'win32' }, async t => {
|
||||
test('Windows: a drive picked as D:\\ becomes D:\\Strata, and the person is not asked for a different folder', { skip: process.platform !== 'win32' && 'drive-letter paths: Windows only' }, async t => {
|
||||
const dir = scratch(t);
|
||||
const letter = 'QRSTUVWXYZ'.split('').find(l => !existsSync(`${l}:\\`));
|
||||
assert.ok(letter, 'a free drive letter');
|
||||
@@ -84,7 +84,7 @@ test('Windows: a drive picked as D:\\ becomes D:\\Strata, and the person is not
|
||||
|
||||
function readdirNames(dir) { return execFileSync('cmd', ['/c', 'dir', '/b', dir]).toString().split(/\r?\n/).filter(Boolean); }
|
||||
|
||||
test('Windows: the model library on a whole drive goes to X:\\Strata\\Models', { skip: process.platform !== 'win32' }, t => {
|
||||
test('Windows: the model library on a whole drive goes to X:\\Strata\\Models', { skip: process.platform !== 'win32' && 'drive-letter paths: Windows only' }, t => {
|
||||
const dir = scratch(t), home = scratch(t);
|
||||
const letter = 'QRSTUVWXYZ'.split('').find(l => !existsSync(`${l}:\\`));
|
||||
execFileSync('subst', [`${letter}:`, dir]);
|
||||
|
||||
@@ -8,7 +8,7 @@ const four = { gpus: [0, 1, 2, 3].map(i => amd(i)) };
|
||||
const old = { arch: 'deepseek41', backend: 'hip', gpu: 0, gpus_asked: true };
|
||||
const linux = process.platform !== 'win32';
|
||||
|
||||
test('setup\'s own one-card pick is not a choice (Linux, several cards)', { skip: !linux }, () => {
|
||||
test('setup\'s own one-card pick is not a choice (Linux, several cards)', { skip: !linux && 'reads the Linux card layout (Linux only)' }, () => {
|
||||
assert.equal(setupPickedOneCard(old, four), true);
|
||||
});
|
||||
test('another card, a marker, roles, a list or another model are the person\'s', () => {
|
||||
|
||||
@@ -13,7 +13,7 @@ test('the UTF-8 line comes first and leaves the script alone', () => {
|
||||
|
||||
// #3202: PowerShell writes to a pipe in the OEM code page (cp866 on a Russian Windows, cp437 elsewhere); a Cyrillic folder came back
|
||||
// as mojibake or "?". The same happens on any non-UTF-8 OEM page, so this fails on this machine without the line, too.
|
||||
test('a Cyrillic temp folder name survives PowerShell through a pipe', { skip: process.platform !== 'win32' }, t => {
|
||||
test('a Cyrillic temp folder name survives PowerShell through a pipe', { skip: process.platform !== 'win32' && 'needs Windows PowerShell (Windows only)' }, t => {
|
||||
const word = String.fromCharCode(0x414, 0x43e, 0x43a, 0x443, 0x43c, 0x435, 0x43d, 0x442, 0x44b); // "Документы"
|
||||
const dir = mkdtempSync(join(tmpdir(), word + '-'));
|
||||
t.after(() => rmSync(dir, { recursive: true, force: true }));
|
||||
|
||||
File diff suppressed because it is too large.
Load diff
@@ -259,7 +259,6 @@ For testing. All are off unless set; the defaults are described above.
|
||||
| `STRATA_METAL_MTP_Q8=0` | The draft layer's matrices in BF16 instead of Q8_0 |
|
||||
| `STRATA_METAL_MTP_CHAIN=0` | One GPU submission per draft step instead of one per round |
|
||||
| `STRATA_METAL_FUSE=0`, `STRATA_METAL_LAZY_FOLD=0` | Turn off two decode fusions (same tokens, slower) |
|
||||
| `STRATA_METAL_PARTIAL=1`, `STRATA_METAL_ROUTE=1` | Two decode variants that measured no faster. `ROUTE` is also wrong on an M1 Max (stale routing logits) |
|
||||
| `STRATA_METAL_PROMPT_EXPERTS=v3` | The M5 layout for the prompt's expert GEMMs. Fails on an M1 Max, so it is off there (default v2) |
|
||||
| `STRATA_METAL_STREAM=1` / `=0` | Stream the experts from the SSD even when the pack fits / never (the default decides by size) |
|
||||
| `STRATA_METAL_CONTIG=0` | Streaming: read the model file even when the experts' copy is there |
|
||||
|
||||
@@ -0,0 +1,364 @@
|
||||
# Engine settings
|
||||
|
||||
<!-- Generated from data/settings-register.json by tools/settings_register.py --docs. Do not edit by hand; change the register. -->
|
||||
|
||||
The engine reads its settings from environment variables whose names start with `STRATA_`. The Terminal and the app choose what
|
||||
your hardware needs, so most people never set one. This page lists the ones meant for people: switches that are off until you
|
||||
turn them on, and switches that take a newer default back to the earlier behaviour. The other settings are for measuring and
|
||||
debugging and are kept in `data/settings-register.json`.
|
||||
|
||||
A number or a word is set like any environment variable, for example `STRATA_KVA=1`. In the Terminal, `strata launch --set NAME=value`
|
||||
puts it in the model's run configuration.
|
||||
|
||||
## Switches you can turn on (280)
|
||||
|
||||
Off unless you set them. The last column says where each one is described.
|
||||
|
||||
| Setting | Value | Applies to | Described in |
|
||||
|---|---|---|---|
|
||||
| `STRATA_ADAPT_FETCH` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ADAPT_FETCH_MIN` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ADAPT_OVERLAP` | text | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_AMD_COMMUNITY_TUNING` | on / off | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_ARENA_FALLBACK_RESIDENT` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ARENA_MMAP` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ARENA_PIN_GIB` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ARENA_RAM_WAIT_S` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ARGMAX_MULTI` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ARGMAX_ROWS_SM70` | number | NVIDIA and AMD | [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_ATTN_PRE75` | on / off | NVIDIA and AMD, Intel Arc | [NVIDIA_V100.md](NVIDIA_V100.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_ATTN_RS` | number | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_AUX_CPUS` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_AUX_STAGER` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_BATCH_DECODE_SHARE` | number | NVIDIA and AMD, Intel Arc | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_BATCH_MTP` | on / off | NVIDIA and AMD, Intel Arc | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_BATCH_YIELD_CHUNKS` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_CACHE_DIR` | on / off | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_CACHE_MESSAGE_BOUNDARY` | number | NVIDIA and AMD, Intel Arc | [MESSAGE_BOUNDARY_CACHE.md](MESSAGE_BOUNDARY_CACHE.md), [PROMPT_CACHE_TAIL.md](PROMPT_CACHE_TAIL.md) |
|
||||
| `STRATA_CKPT_ASYNC` | on / off | NVIDIA | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_CKPT_REREAD` | on / off | NVIDIA and AMD, Intel Arc | [KVA.md](KVA.md), [PROMPT_CACHE_TAIL.md](PROMPT_CACHE_TAIL.md) |
|
||||
| `STRATA_CKPT_THREADS` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_CKPT_VRAM_SLOTS` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_COLD_OVERLAP` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.0.md](release-notes-0.2.0.md) |
|
||||
| `STRATA_COLD_TRACE` | on / off | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_COPY_ENGINE` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_COPY_HIP_PRIO` | number | NVIDIA and AMD | - |
|
||||
| `STRATA_DBG_GDN` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_DENSE_F32` | number | NVIDIA | [COMMUNITY_BENCHMARKS.md](COMMUNITY_BENCHMARKS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_DENSE_MMQ` | on / off | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_DEQUANT_SIDE` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_DISJOINT_ADAPT` | number | NVIDIA and AMD, Intel Arc | [DISJOINT_EXPERT_CACHE.md](DISJOINT_EXPERT_CACHE.md) |
|
||||
| `STRATA_DMA_BATCH` | number | NVIDIA and AMD | [BATCHED_DMA.md](BATCHED_DMA.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_DOWN_WS` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_DS4_READ_PIECES` | number | NVIDIA and AMD | - |
|
||||
| `STRATA_DS4_SLOT_GB` | number | NVIDIA and AMD | [DEEPSEEK41.md](DEEPSEEK41.md) |
|
||||
| `STRATA_DS4_SLOT_KEEP` | number | NVIDIA and AMD | [DEEPSEEK41.md](DEEPSEEK41.md) |
|
||||
| `STRATA_DUMP_FIRST_LOGITS` | on / off | NVIDIA and AMD, Intel Arc | [DGX_SPARK_LINUX.md](DGX_SPARK_LINUX.md) |
|
||||
| `STRATA_DUMP_MOE_INPUT` | number | NVIDIA and AMD | [NVFP4.md](NVFP4.md) |
|
||||
| `STRATA_EMBED_VRAM` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_EXCHANGE_ROTATE` | on / off | NVIDIA and AMD, Intel Arc, CPU | [DETAILS.md](DETAILS.md), [EXCHANGE_ROTATION.md](EXCHANGE_ROTATION.md) |
|
||||
| `STRATA_EXPDEC_MINB` | number | NVIDIA and AMD | [EXPERTDEC.md](EXPERTDEC.md) |
|
||||
| `STRATA_EXPDEC_MINBLK` | number | NVIDIA and AMD | [EXPERTDEC.md](EXPERTDEC.md) |
|
||||
| `STRATA_EXPDEC_RB` | number | NVIDIA and AMD | [EXPERTDEC.md](EXPERTDEC.md) |
|
||||
| `STRATA_EXPERT_NIBBLE` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_EXPERT_NIBBLE_FMTS` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_EXPERT_NIBBLE_RECHECK` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_EXPERT_NIBBLE_SKIP` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_EXPERT_V2K` | on / off | NVIDIA and AMD, Intel Arc | [QUANTS.md](QUANTS.md), [QWEN36.md](QWEN36.md) |
|
||||
| `STRATA_EXP_MODE` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [NVIDIA_V100.md](NVIDIA_V100.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_F16_RANGE` | on / off | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_FETCH_THREADS` | number | NVIDIA and AMD, Intel Arc | [UNSLOTH_Q4.md](UNSLOTH_Q4.md) |
|
||||
| `STRATA_FILE_CACHE_OLD_RULE` | on / off | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_FILE_RELEASE` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_FILL_AHEAD` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_FS_SLOTS` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_GDN_CHUNK` | number | NVIDIA and AMD | [GDN_CONTRACT.md](GDN_CONTRACT.md), [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_GDN_PIPELINE` | number | NVIDIA and AMD, Intel Arc | [GDN_CONTRACT.md](GDN_CONTRACT.md) |
|
||||
| `STRATA_GDN_REC_HEADS` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_GDN_TC` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [GDN_CONTRACT.md](GDN_CONTRACT.md) |
|
||||
| `STRATA_GFX12_CHUNK` | text | AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GFX12_RING` | on / off | AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GFX906_ATTN_QUERY_SWIZZLE` | on / off | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_GFX906_ATTN_REDUCE12` | on / off | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_GRAPH_PRIO` | number | NVIDIA | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GR_DOWN_TG` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GR_SPLIT` | on / off | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HANDOFF_PIPE` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_HC_PERSIST` | number | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HC_Q8` | on / off | NVIDIA and AMD, Intel Arc, Mac | [AMD_HIP.md](AMD_HIP.md), [MAC.md](MAC.md), [QUANTS.md](QUANTS.md) |
|
||||
| `STRATA_HC_REQ8` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_HEAD_HOST` | on / off | NVIDIA and AMD | - |
|
||||
| `STRATA_HIPBLASLT_TABLE` | on / off | AMD | [AMD_HIP.md](AMD_HIP.md), [release-notes-0.2.1.md](release-notes-0.2.1.md) |
|
||||
| `STRATA_HIPBLASLT_TUNING` | on / off | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [AMD_HIP_PERFORMANCE.md](AMD_HIP_PERFORMANCE.md), [RADEON_780M.md](RADEON_780M.md) |
|
||||
| `STRATA_HIPBLASLT_VERBOSE` | text | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HIP_ADAPT_KERNEL_COPY` | on / off | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_HIP_PROMPT_F16` | on / off | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HIP_ROUTER_OLD` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HIP_WMMA_MIN_Q` | number | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_HOST_CORE` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_I8_DBG` | number | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_I8_GRF` | number | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_I8_MT` | number | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_I8_WG` | number | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_IO_PF_AHEAD` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IO_PF_GATE` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.0.md](release-notes-0.2.0.md) |
|
||||
| `STRATA_IO_PF_GATE_LO` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IO_PF_GATE_USEFUL` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IO_PF_STAGE` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IO_PF_THREADS` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IO_PREFETCH` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md), [RADEON_780M.md](RADEON_780M.md) |
|
||||
| `STRATA_IO_PREFETCH_DEPTH` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IQ3S_MT1` | on / off | CPU | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IQ3S_WIDE` | number | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_IQ4_W4A8` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_IQ512_FUSED_MAX` | number | CPU | - |
|
||||
| `STRATA_IQ_STAGE_GRID` | on / off | NVIDIA and AMD, Intel Arc | [GYRO.md](GYRO.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_IQ_STAGE_TABLES` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_KVA` | on / off | NVIDIA and AMD | [KVA.md](KVA.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_KVA_ANY_MODEL` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_CAPTURE` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_CAPTURE_FP32` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_GDN_NORM` | number | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_MIN_PROMPT` | number | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_MODE` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_ONLY` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_ORACLE` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_SKIP` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KVA_SUFFIX` | number | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_KV_GROW` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_KV_GROW_HOLD` | number | NVIDIA and AMD | - |
|
||||
| `STRATA_KV_HOST_DMA` | on / off | AMD, Intel Arc | [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_KV_HOST_REGISTER` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_KV_PREFETCH` | number | NVIDIA and AMD, Intel Arc | [KV_PREFETCH.md](KV_PREFETCH.md) |
|
||||
| `STRATA_KV_ROT` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md), [SPLIT_SERVE.md](SPLIT_SERVE.md) |
|
||||
| `STRATA_LEND_MIRROR` | number | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_LM_HIP` | on / off | AMD | - |
|
||||
| `STRATA_LOGITS_DUMP` | on / off | NVIDIA and AMD | [KVA.md](KVA.md) |
|
||||
| `STRATA_LOGPOS_TOPK` | number | NVIDIA and AMD, Intel Arc | [QUANTS.md](QUANTS.md), [UNSLOTH_Q4.md](UNSLOTH_Q4.md) |
|
||||
| `STRATA_LOOKUP_MAX_T` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_MEMORY_LIMIT_MIB` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_METAL_ARENA_GB` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_CONTIG` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_F16` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_GYRO_TILES` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_KEEP_FREE_GB` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_KEEP_HOST_WEIGHTS` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_KPROF` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_LEAVE_SHARE` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_MTP_CHAIN` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_PROMPT` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_PROMPT_CHUNK` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_PROMPT_EXPERTS` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_STREAM` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_TRACE` | on / off | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_METAL_WIRE_ARENA` | number | Mac | [MAC.md](MAC.md) |
|
||||
| `STRATA_MIRROR_MIB` | number | Intel Arc | [INTEL.md](INTEL.md), [INTEL_ARC.md](INTEL_ARC.md) |
|
||||
| `STRATA_MMQ_DIRECT` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_MMQ_GROUP_SORT` | text | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_MMQ_WS` | number | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_MMVQ_IL_89` | text | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [MMVQ_IL_TABLE.md](MMVQ_IL_TABLE.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_MMVQ_IL_PF` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [MMVQ_IL_TABLE.md](MMVQ_IL_TABLE.md) |
|
||||
| `STRATA_MMVQ_IL_ROWS` | on / off | NVIDIA and AMD | [MMVQ_IL_TABLE.md](MMVQ_IL_TABLE.md) |
|
||||
| `STRATA_MMVQ_WIDE_32` | number | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_MTP_KEEP` | on / off | NVIDIA and AMD | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_MTP_Q8_SM60` | on / off | NVIDIA and AMD | [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_MTP_ROOM_MIN` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_NIB_GU_PROTO` | number | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_NO_ARENA_THP` | on / off | NVIDIA and AMD, Intel Arc | [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_NO_AVXVNNI` | number | CPU | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_NO_TQ_KERNELS` | number | CPU | [GYRO.md](GYRO.md) |
|
||||
| `STRATA_OWNED_PRICE` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PARK_ASYNC` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PARK_TRACE` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PARK_VRAM_MIB` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PARTIAL_PIN` | number | NVIDIA and AMD, Intel Arc | [UNSLOTH_Q4.md](UNSLOTH_Q4.md) |
|
||||
| `STRATA_PARTIAL_PIN_GIB` | number | NVIDIA and AMD, Intel Arc | [UNSLOTH_Q4.md](UNSLOTH_Q4.md) |
|
||||
| `STRATA_PATTN_SM60` | number | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_PA_V2T` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PCIE_AUTO_CANON` | on / off | NVIDIA and AMD | - |
|
||||
| `STRATA_PCIE_FRAC_DEFAULT` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PCIE_K_RULE` | text | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PCIE_MIN1` | number | NVIDIA and AMD | - |
|
||||
| `STRATA_PEER_NO_P2P` | number | NVIDIA and AMD | [SECOND_GPU.md](SECOND_GPU.md) |
|
||||
| `STRATA_PENALTY_PROMPT` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_BATCH` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_DMA_PAGEABLE` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_PF_FUSED_F16` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_FUSED_IQ2S` | text | NVIDIA and AMD | - |
|
||||
| `STRATA_PF_FUSED_KQ` | on / off | NVIDIA and AMD, Intel Arc | [QUANTS.md](QUANTS.md), [STRIX_HALO.md](STRIX_HALO.md) |
|
||||
| `STRATA_PF_FUSED_NATIVE` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_FUSED_NT` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_FUSED_Q2N` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_PF_FUSED_SM75` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_FUSED_WS` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_GROUP_SCRATCH` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_HC_REGION` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_I8` | on / off | Intel Arc | [INTEL_INT8_EXPERTS.md](INTEL_INT8_EXPERTS.md) |
|
||||
| `STRATA_PF_PEER_CPU` | number | NVIDIA and AMD | - |
|
||||
| `STRATA_PF_SWITCH_MIN_T` | number | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [STRIX_HALO.md](STRIX_HALO.md) |
|
||||
| `STRATA_PF_WMMA` | on / off | NVIDIA and AMD | [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_PF_WMMA2` | on / off | NVIDIA and AMD | [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_PIN_BY_NEED` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [DOCKER.md](DOCKER.md), [TERMINAL_LAUNCH.md](TERMINAL_LAUNCH.md) |
|
||||
| `STRATA_PIN_GUARD` | number | NVIDIA and AMD | [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_PIN_RESERVE_GIB` | number | NVIDIA and AMD | [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_PIPED_CONCURRENT` | text | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md), [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PIPED_CONCURRENT_PRIO` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PIPED_DECODE_MAX_MS` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PIPED_DECODE_MIN_MS` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PIPED_MTP` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_PIPED_READ_PIECE` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PIPELINE_AUTO` | text | NVIDIA and AMD, Intel Arc | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_PIPELINE_DEBUG` | on / off | NVIDIA and AMD, Intel Arc | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_PIPE_ADAPT_EVERY` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_PIPE_BYIELD_NODRAIN` | on / off | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_PLE_COMPACT` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PLE_COPY_STREAM` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_POOL_CORES` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_PREFILL_AUTO_MAX` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_BORROW` | text | NVIDIA, AMD | [STRIX_HALO.md](STRIX_HALO.md), [STRIX_HALO_WINDOWS.md](STRIX_HALO_WINDOWS.md) |
|
||||
| `STRATA_PREFILL_CPU_SHARE_BATCH` | on / off | NVIDIA and AMD | - |
|
||||
| `STRATA_PREFILL_CPU_SHARE_MAX` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_EQUAL` | on / off | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_FIRST` | number | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_PREFILL_HELP` | number | NVIDIA and AMD, Intel Arc | [BATCHING.md](BATCHING.md), [GYRO.md](GYRO.md), [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_PREFILL_HELP_FRAC` | number | NVIDIA and AMD, Intel Arc | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_PREFILL_HOST_PLAN` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.0.md](release-notes-0.2.0.md) |
|
||||
| `STRATA_PREFILL_ISSUER` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_LAYER_MAJOR` | on / off | NVIDIA and AMD | [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_PREFILL_MMQ` | number | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [AMD_HIP_PERFORMANCE.md](AMD_HIP_PERFORMANCE.md), [GYRO.md](GYRO.md) |
|
||||
| `STRATA_PREFILL_PIPE` | number | NVIDIA and AMD | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_PREFILL_ROUTED_STREAM` | text | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_ROUTED_STREAM_DEPTH` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PREFILL_STREAM_ALL` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_PREFILL_STREAM_MIN` | number | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [STRIX_HALO.md](STRIX_HALO.md) |
|
||||
| `STRATA_PRIMARY_DEVICE` | number | NVIDIA and AMD | [SECOND_GPU.md](SECOND_GPU.md) |
|
||||
| `STRATA_PROMPT_ATTN_IMMA` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PROMPT_DIRECT` | on / off | NVIDIA and AMD | [QUANTS.md](QUANTS.md) |
|
||||
| `STRATA_PW_WITH_BATCH` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_Q2_BITPLANE` | on / off | CPU | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_Q36_FIXED_SLOTS` | on / off | NVIDIA and AMD | [QWEN36.md](QWEN36.md) |
|
||||
| `STRATA_Q36_READ_WINDOWS` | number | NVIDIA and AMD | [QWEN36.md](QWEN36.md) |
|
||||
| `STRATA_Q36_VERIFY` | on / off | NVIDIA and AMD | [QWEN36.md](QWEN36.md) |
|
||||
| `STRATA_Q6_PLANAR` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_Q8_GATHER_WIDE` | on / off | NVIDIA | [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_Q8_SM60` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_QSA_CLUSTER` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_QSA_SPLIT_ATTN` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_RAM_ADAPT` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_RAM_KEEP_GIB` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_RDNA2_SGEMM` | number | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_RESIDENT_HEADROOM_GIB` | number | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md), [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_ROCBLAS_TUNING` | on / off | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_ROCBLAS_VERBOSE` | text | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_ROUTE_PER_TOKEN` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_ROUTE_PRIOR` | text | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.0.md](release-notes-0.2.0.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_ROUTE_RESIDENT` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ROUTE_RESIDENT_MTP` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SCORES_HALVE` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SCORES_THREAD_TRACE` | text | AMD | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_SELECT_LONG` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_SELECT_LONG_AT` | number | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_SELECT_SIMT` | on / off | NVIDIA and AMD | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_SELECT_TWO_PASS_CHECK` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SELECT_TWO_PASS_MIN` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SM120_PERF` | number | NVIDIA | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SM121_PERF` | number | NVIDIA and AMD | [DGX_SPARK_LINUX.md](DGX_SPARK_LINUX.md) |
|
||||
| `STRATA_SM70_TABLE` | number | NVIDIA and AMD | [COMMUNITY_BENCHMARKS.md](COMMUNITY_BENCHMARKS.md), [NVIDIA_V100.md](NVIDIA_V100.md) |
|
||||
| `STRATA_SM75_PERF` | number | NVIDIA | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SPEC_COUPLED` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SPEC_GUMBEL` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SPEC_MIN_TEMP` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SPEC_PROB` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_SPEC_PROB_GATE` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SPEC_WATCH` | number | NVIDIA and AMD | [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_SPLIT_HEADROOM` | number | Intel Arc | [INTEL_ARC.md](INTEL_ARC.md) |
|
||||
| `STRATA_SPLIT_HEADROOM_TOL` | number | Intel Arc | [INTEL_ARC.md](INTEL_ARC.md) |
|
||||
| `STRATA_SPLIT_MTP_BATCH` | text | NVIDIA and AMD | [BATCHING.md](BATCHING.md), [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_SPLIT_OVERLAP` | number | NVIDIA and AMD | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_SPLIT_UNIFORM` | on / off | NVIDIA and AMD | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_SPLIT_WAIT_MS` | number | NVIDIA and AMD | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_SSD_GBPS` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_STAGER_SLEEP` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_STAGER_SSD` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_STAGE_PIN` | number | NVIDIA and AMD | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_STATE_HASH` | on / off | NVIDIA and AMD, Intel Arc | [SPLIT_SERVE.md](SPLIT_SERVE.md) |
|
||||
| `STRATA_SYCL_GEMM_STAGE` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_TEST_BELL_SLEEP_US` | number | NVIDIA and AMD | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_TOPK_MULTI` | on / off | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_TOPK_OLD` | on / off | NVIDIA and AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_TQ_ISA` | on / off | CPU | [GYRO.md](GYRO.md) |
|
||||
| `STRATA_UMA_HEADROOM_GIB` | number | NVIDIA and AMD, Intel Arc | [STRIX_HALO.md](STRIX_HALO.md) |
|
||||
| `STRATA_UNBUFFERED_LOAD` | on / off | NVIDIA and AMD, Intel Arc | [QUANTS.md](QUANTS.md), [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_VERIFY_ALL_SLOTS` | on / off | Intel Arc | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_VERIFY_EARLY_PLE_FLAG` | number | NVIDIA and AMD | [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_VERIFY_NO_HOST` | on / off | Intel Arc | [A770_SPEED_SPEC.md](A770_SPEED_SPEC.md), [INTEL.md](INTEL.md), [INTEL_ARC.md](INTEL_ARC.md) |
|
||||
| `STRATA_VRAM_FREE` | on / off | NVIDIA and AMD | [DETAILS.md](DETAILS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_WATCHDOG_S` | number | NVIDIA and AMD, Intel Arc | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_WDDM_BUDGET` | number | AMD, Intel Arc | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_WS_FAST` | number | NVIDIA and AMD | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
|
||||
## Switches that go back to earlier behaviour (58)
|
||||
|
||||
These behaviours are on by default. Setting the switch returns the earlier behaviour; the notes in the register say which
|
||||
version made it the default.
|
||||
|
||||
| Setting | Value | Applies to | Named in the engine since | Described in |
|
||||
|---|---|---|---|---|
|
||||
| `STRATA_ADAPT_SIGNAL` | on / off | NVIDIA and AMD | v0.1.43.1 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_ATTN_V3` | on / off | NVIDIA and AMD | v0.1.45 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_BATCH_GROUPS_AUTO` | number | Intel Arc | v0.2.9 | - |
|
||||
| `STRATA_BATCH_GROUPS_NV_SERIAL` | text | NVIDIA and AMD | v0.2.1 | [BATCHING.md](BATCHING.md), [release-notes-0.2.1.md](release-notes-0.2.1.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_BATCH_MTP_SERIAL` | on / off | NVIDIA and AMD | v0.1.43.1 | [BATCHING.md](BATCHING.md) |
|
||||
| `STRATA_DEFAULTS` | text | NVIDIA and AMD | v0.2.7 | [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md), [DGX_SPARK_LINUX.md](DGX_SPARK_LINUX.md) |
|
||||
| `STRATA_DEQUANT_WIDE` | number | NVIDIA and AMD | v0.1.44 | [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_FILE_TIER_IDLE` | number | NVIDIA and AMD | v0.2.12 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GDN_CHUNKED` | number | NVIDIA and AMD | v0.1.41 | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md), [GDN_CONTRACT.md](GDN_CONTRACT.md) |
|
||||
| `STRATA_GFX1151_CHUNK` | text | AMD | v0.2.3 | [release-notes-0.2.3.md](release-notes-0.2.3.md), [STRIX_HALO.md](STRIX_HALO.md), [STRIX_HALO_WINDOWS.md](STRIX_HALO_WINDOWS.md) |
|
||||
| `STRATA_GFX1151_DEFAULTS` | on / off | NVIDIA and AMD | v0.1.40.1 | [STRIX_HALO.md](STRIX_HALO.md), [STRIX_HALO_WINDOWS.md](STRIX_HALO_WINDOWS.md) |
|
||||
| `STRATA_GFX12_DEFAULTS` | on / off | NVIDIA and AMD | v0.1.42 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_GR_V3` | number | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_HC_Q8_ALL_STAGES` | on / off | NVIDIA and AMD | v0.2.12 | - |
|
||||
| `STRATA_HIP_QUEUE_PRIME` | on / off | NVIDIA and AMD | v0.2.1 | [AMD_HIP.md](AMD_HIP.md), [release-notes-0.2.1.md](release-notes-0.2.1.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_IGPU_DEFAULTS` | on / off | NVIDIA and AMD | v0.2.1 | [AMD_HIP.md](AMD_HIP.md) |
|
||||
| `STRATA_IGPU_PROMPT_SET` | on / off | NVIDIA and AMD | v0.2.1 | [AMD_HIP.md](AMD_HIP.md), [RADEON_780M.md](RADEON_780M.md), [release-notes-0.2.1.md](release-notes-0.2.1.md) |
|
||||
| `STRATA_IL_PAD` | number | NVIDIA and AMD | v0.2.10 | - |
|
||||
| `STRATA_IQ_MT_MIN` | on / off | NVIDIA and AMD, CPU | v0.1.40.1 | [BATCHING.md](BATCHING.md), [DETAILS.md](DETAILS.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_KEEP_MAPPING` | on / off | NVIDIA and AMD | v0.1.41 | - |
|
||||
| `STRATA_METAL_LDB` | on / off | Mac | v0.2.10 | - |
|
||||
| `STRATA_METAL_MTP_OVERLAP` | on / off | Mac | v0.2.10 | - |
|
||||
| `STRATA_METAL_MTP_Q8` | on / off | Mac | v0.1.44 | [MAC.md](MAC.md) |
|
||||
| `STRATA_MMVQ_A2` | number | Intel Arc | v0.1.40.1 | [INTEL.md](INTEL.md) |
|
||||
| `STRATA_MMVQ_IL` | on / off | NVIDIA and AMD | v0.1.40.2 | [MMVQ_IL_TABLE.md](MMVQ_IL_TABLE.md), [NVIDIA_V100.md](NVIDIA_V100.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_MMVQ_IL_SM80` | text | NVIDIA and AMD | v0.2.7 | [DETAILS.md](DETAILS.md), [MMVQ_IL_TABLE.md](MMVQ_IL_TABLE.md), [release-notes-0.2.7.md](release-notes-0.2.7.md) |
|
||||
| `STRATA_MTP_BATCH_RING` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | - |
|
||||
| `STRATA_NAN_GUARD` | number | NVIDIA and AMD | v0.1.43 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_NO_IQ512` | on / off | CPU | v0.1.40.1 | - |
|
||||
| `STRATA_NUMA_POOL` | on / off | CPU | v0.2.2 | [DETAILS.md](DETAILS.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_NUMA_SPLIT` | text | CPU | v0.2.2 | [DETAILS.md](DETAILS.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_NV_CHUNK` | on / off | NVIDIA | v0.2.11 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PF_W16_CACHE_MIB` | on / off | NVIDIA and AMD | v0.2.9 | - |
|
||||
| `STRATA_PF_XMX` | text | Intel Arc | v0.1.42 | [INTEL.md](INTEL.md), [release-notes-0.2.0.md](release-notes-0.2.0.md) |
|
||||
| `STRATA_PIPELINE_ADAPT_ASYNC` | number | NVIDIA and AMD | v0.1.40.2 | - |
|
||||
| `STRATA_PLE_AUTO_MMAP` | number | NVIDIA and AMD, Intel Arc | v0.1.43.1 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_PROMPT_ATTN_ESIMD` | number | Intel Arc | v0.2.9 | - |
|
||||
| `STRATA_PROMPT_ATTN_OLD` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [NVIDIA_V100.md](NVIDIA_V100.md), [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_Q2_WARP` | on / off | NVIDIA and AMD | v0.1.45 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_Q36_CHUNK` | number | NVIDIA and AMD | v0.1.45 | [QWEN36.md](QWEN36.md) |
|
||||
| `STRATA_QFUSE_HC_LOCAL` | text | NVIDIA and AMD | v0.1.43 | - |
|
||||
| `STRATA_QSA_EARLY_EXIT` | number | NVIDIA and AMD, Intel Arc | v0.1.40.1 | - |
|
||||
| `STRATA_QSA_FUSE_NR` | on / off | NVIDIA and AMD | v0.2.10 | - |
|
||||
| `STRATA_REMOTE_SPIN` | on / off | NVIDIA, AMD, Intel Arc | v0.1.40.1 | - |
|
||||
| `STRATA_RESIDENT_PIN` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [DETAILS.md](DETAILS.md), [TROUBLESHOOTING.md](TROUBLESHOOTING.md), [UNSLOTH_Q4.md](UNSLOTH_Q4.md) |
|
||||
| `STRATA_RING_BYTES` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_S2_SWIGLU_Q8` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | - |
|
||||
| `STRATA_SCORES_THREAD` | number | NVIDIA and AMD | v0.1.45 | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_SELECT_OLD` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | - |
|
||||
| `STRATA_SM60_TABLE` | number | NVIDIA and AMD | v0.1.43 | [OLDER_GPUS.md](OLDER_GPUS.md) |
|
||||
| `STRATA_SPLIT_CHUNK16K` | number | NVIDIA and AMD | v0.1.44 | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_SPLIT_OWN` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [MULTI_GPU.md](MULTI_GPU.md) |
|
||||
| `STRATA_TOPK_SAMPLE` | number | NVIDIA and AMD | v0.1.45 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_TOPK_SPLIT` | number | NVIDIA and AMD | v0.1.45 | [DETAILS.md](DETAILS.md) |
|
||||
| `STRATA_TOPK_STREAM` | on / off | NVIDIA and AMD, Intel Arc | v0.1.40.1 | [AMD_HIP.md](AMD_HIP.md), [DETAILS.md](DETAILS.md), [TROUBLESHOOTING.md](TROUBLESHOOTING.md) |
|
||||
| `STRATA_TQ_STAGE` | on / off | NVIDIA and AMD | v0.2.2 | [GYRO.md](GYRO.md), [release-notes-0.2.2.md](release-notes-0.2.2.md) |
|
||||
| `STRATA_UMA_RESERVE_GIB` | on / off | NVIDIA and AMD | v0.2.2 | [release-notes-0.2.2.md](release-notes-0.2.2.md), [STRIX_HALO.md](STRIX_HALO.md) |
|
||||
| `STRATA_XN_F16` | on / off | NVIDIA and AMD | v0.1.45 | [DETAILS.md](DETAILS.md) |
|
||||
@@ -6,8 +6,8 @@
|
||||
//
|
||||
// #3702: on an unequal pair (RTX 3060 + RTX 5060 Ti) the last card was left with 68 MiB after loading although
|
||||
// 700 + 96 MiB were reserved for it, and the verify windows' graphs (captured on first use, 4-6 MiB each and once 42 MiB)
|
||||
// then ran it out of memory. `proposed_window_mib` below is a PROPOSAL for a larger window term; nothing calls it
|
||||
// (it needs the model gate in phaseA-tests/gh027/gates/3702-later-reserve.sh before it can become the default).
|
||||
// then ran it out of memory. A larger window term was proposed for it; the proposal is not part of the default placement
|
||||
// and lives in research/stage-reserve-proposal/ (nothing in the engine calls it).
|
||||
#pragma once
|
||||
|
||||
#include <algorithm>
|
||||
@@ -38,24 +38,4 @@ inline int64_t room_bytes(int64_t free_bytes, const Inputs& in) {
|
||||
return std::max<int64_t>(free_bytes - (reserve_mib(in) << 20), 0);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------------------------------- the proposal
|
||||
constexpr int64_t kGraphMib = 6; ///< one window graph, captured on first use (log of #3702: 4, 6, 6)
|
||||
constexpr int64_t kInstantiateBurstMib = 48; ///< the largest single step seen (42 MiB in #3702): room kept for one more
|
||||
|
||||
/// PROPOSAL, unused: the window term of a later stage when the windows of every size 1..max_window_t can still be
|
||||
/// captured after loading: the first window's arena (kWindowMib already holds it), a graph for each further size, and
|
||||
/// one burst.
|
||||
inline int64_t proposed_window_mib(int max_window_t) {
|
||||
const int64_t sizes = std::max<int64_t>(max_window_t, 1);
|
||||
return kWindowMib + (sizes - 1) * kGraphMib + kInstantiateBurstMib;
|
||||
}
|
||||
|
||||
/// PROPOSAL, unused: reserve_mib with the proposed window term on a later stage.
|
||||
inline int64_t proposed_reserve_mib(const Inputs& in, int max_window_t) {
|
||||
Inputs p = in;
|
||||
int64_t r = reserve_mib(p);
|
||||
if (in.later) r += proposed_window_mib(max_window_t) - kWindowMib;
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace strata::program::stage_reserve
|
||||
@@ -70,58 +70,6 @@ inline int choose_window(const WindowRuleIn& in) {
|
||||
return T;
|
||||
}
|
||||
|
||||
/// STRATA_DF2_GATE=1 (opt-in, measured on the Mac): DFlash2's window from a calibrated survival instead of a fixed confidence. The selector's cumulative confidence C_j of
|
||||
/// draft j over-states the real chance that drafts 1..j all hold (M5 Max UD-Q4_K_XL: C = 0.84, survival 0.70 at j = 5), and the right confidence differs by request
|
||||
/// (a story wants a narrower window than JSON). Per request this keeps, for every position j and confidence bucket, how often the first j drafts held, shrunk toward a
|
||||
/// prior of C^1.3, and extends the window to row j+1 while survival * (ms per committed token) >= row_ms (what one more verify row costs).
|
||||
struct AdaptiveGate {
|
||||
static constexpr int NB = 6;
|
||||
double n[8][NB] = {}, h[8][NB] = {};
|
||||
double ms_per_tok = 0.0; // EMA of the round's wall ms per committed token
|
||||
double row_ms = 3.0; // one more verify row (M5 Max: 2.65-3.3 ms)
|
||||
double prior_n = 3.0;
|
||||
static int bucket(double c) { return c < 0.15 ? 0 : c < 0.3 ? 1 : c < 0.5 ? 2 : c < 0.7 ? 3 : c < 0.85 ? 4 : 5; }
|
||||
double survival(int j, double c) const {
|
||||
const int b = bucket(c);
|
||||
const double prior = std::pow(c, 1.3);
|
||||
return (h[j][b] + prior_n * prior) / (n[j][b] + prior_n);
|
||||
}
|
||||
/// rows of the next window: 1 + the drafts kept (cap = widest window)
|
||||
int choose(const float* dprob, int cap) const {
|
||||
int T = 1;
|
||||
double chain = 1.0;
|
||||
const double v = ms_per_tok > 0.0 ? ms_per_tok : 9.0;
|
||||
while (T < cap) {
|
||||
chain *= (double) dprob[T - 1];
|
||||
if (survival(T, chain) * v < row_ms) break;
|
||||
++T;
|
||||
}
|
||||
return T;
|
||||
}
|
||||
/// after a DFlash2 round of T rows, a drafts accepted: drafts 1..T-1 were offered, draft j survived when a >= j
|
||||
void observe(const float* dprob, int T, int a) {
|
||||
double chain = 1.0;
|
||||
for (int j = 1; j < T && j < 8; ++j) {
|
||||
chain *= (double) dprob[j - 1];
|
||||
const int b = bucket(chain);
|
||||
n[j][b] += 1.0;
|
||||
h[j][b] += a >= j ? 1.0 : 0.0;
|
||||
}
|
||||
}
|
||||
void observe_round(double round_ms, int committed) {
|
||||
if (committed < 1) return;
|
||||
const double m = round_ms / committed;
|
||||
ms_per_tok = ms_per_tok > 0.0 ? 0.9 * ms_per_tok + 0.1 * m : m;
|
||||
}
|
||||
};
|
||||
|
||||
inline bool df2_gate_on() {
|
||||
static const bool v = [] { const char* e = std::getenv("STRATA_DF2_GATE"); return e != nullptr && e[0] == '1'; }();
|
||||
return v;
|
||||
}
|
||||
inline double df2_gate_row_ms() {
|
||||
static const double v = [] { const char* e = std::getenv("STRATA_DF2_ROW_MS"); return e != nullptr && e[0] != '\0' ? std::atof(e) : 3.0; }();
|
||||
return v;
|
||||
}
|
||||
// (An opt-in calibrated DFlash2 window gate, STRATA_DF2_GATE / STRATA_DF2_ROW_MS, was removed in 0.2.13: it measured 2.7-4.9% slower.)
|
||||
|
||||
} // namespace strata::spec
|
||||
@@ -22,16 +22,11 @@ Exit codes:
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import http.client
|
||||
import json
|
||||
import os
|
||||
import re
|
||||
import shutil
|
||||
import signal
|
||||
import subprocess
|
||||
import sys
|
||||
import tempfile
|
||||
import time
|
||||
from pathlib import Path
|
||||
|
||||
|
||||
@@ -44,7 +44,6 @@ import json
|
||||
import os
|
||||
import re
|
||||
import shutil
|
||||
import signal
|
||||
import socket
|
||||
import struct
|
||||
import subprocess
|
||||
|
||||
@@ -277,12 +277,3 @@ template [[host_name("gr_up_q8")]] kernel decltype(gr_up_f<T_Q8_0>) gr_up_f<T_Q8
|
||||
template [[host_name("gr_downp_q8_8_8")]] kernel decltype(gr_down_f<8, 8, T_Q8_0, T_Q8_0, 1>) gr_down_f<8, 8, T_Q8_0, T_Q8_0, 1>;
|
||||
template [[host_name("gr_downp_q8f_8_8")]] kernel decltype(gr_down_f<8, 8, T_Q8_0, T_F32, 1>) gr_down_f<8, 8, T_Q8_0, T_F32, 1>;
|
||||
template [[host_name("gr_upp_q8")]] kernel decltype(gr_up_f<T_Q8_0, 1>) gr_up_f<T_Q8_0, 1>;
|
||||
|
||||
// wpf: a weight prefetch. Reads n16 uint4 of W (grid-strided) and writes nothing unless the xor of everything is one value in 2^32; the point is only to
|
||||
// pull the next kernel's first weights into the GPU's caches while the small hyper-connection kernels run (STRATA_METAL_WPF_MIB).
|
||||
kernel void wpf(device const uint4* W [[buffer(0)]], device uint* flag [[buffer(1)]], constant uint& n16 [[buffer(2)]],
|
||||
uint gid [[thread_position_in_grid]], uint gsz [[threads_per_grid]]) {
|
||||
uint4 a = 0;
|
||||
for (uint i = gid; i < n16; i += gsz) a ^= W[i];
|
||||
if ((a.x ^ a.y ^ a.z ^ a.w) == 0x7fffffffu) flag[0] = 1;
|
||||
}
|
||||
+12
-72
@@ -430,12 +430,6 @@ private:
|
||||
// the start of its command buffer (one submit and wait less a round). Nothing else reads the GDN state between
|
||||
// the two (the drafter reads the window's rows and its own caches). STRATA_METAL_LAZY_FOLD=0: at commit.
|
||||
int pend_fold_ = 0;
|
||||
// gdn_step's partial commit (mac-kern 56526aed, commit = 0x100 | k): a verify window of T > 1 commits its first
|
||||
// row in place (it is always kept), so commit(1) needs no fold and commit(n) folds rows 1..n-1 only; a window
|
||||
// without the head (the prompt) commits all its rows in the step. Opt-in (STRATA_METAL_PARTIAL=1): at --spec 2
|
||||
// it measured 0.3-0.8% slower on the M5 Max (the in-place state write costs every window, the fold it saves
|
||||
// only the rounds with no accepted draft). Off: the window commits nothing and commit(n) folds n rows.
|
||||
bool partial_ = std::getenv("STRATA_METAL_PARTIAL") && std::string(std::getenv("STRATA_METAL_PARTIAL")) == "1";
|
||||
bool win_head_ = true; // the window being encoded has the head
|
||||
bool last_head_ = true; // the last window had it
|
||||
// token_chain's --fuse (mac-kern): the GDN step writes the prepared y planes itself (no xprep), and a window whose
|
||||
@@ -445,25 +439,9 @@ private:
|
||||
// originals wait one memory latency per loop iteration). Bitwise the same outputs. STRATA_METAL_LDB=0: the original kernels.
|
||||
bool ldb_ = !(std::getenv("STRATA_METAL_LDB") && std::string(std::getenv("STRATA_METAL_LDB")) == "0");
|
||||
const char* k_norm() const { return ldb_ ? "gr_norm2" : "gr_norm"; }
|
||||
// STRATA_METAL_WPF_MIB=N (opt-in screen): while a hyper-connection read's last kernel runs, a prefetch dispatch (no barrier before it) reads the first N MiB of each
|
||||
// weight the next kernels use (wpf_next_), so those kernels find them in the GPU's caches. Bitwise: it changes no value.
|
||||
std::vector<std::pair<uint64_t, uint64_t>> wpf_next_; // (W_ offset, bytes)
|
||||
mtl::Buffer wpf_flag_;
|
||||
size_t wpf_bytes_ = [] { const char* e = std::getenv("STRATA_METAL_WPF_MIB"); return e ? (size_t) std::atoi(e) * 1048576 : (size_t) 0; }();
|
||||
void wpf_add(const std::string& name) {
|
||||
if (!wpf_bytes_) return;
|
||||
const gguf::Tensor* t = f_.find(name);
|
||||
if (t) wpf_next_.push_back({t->offset, std::min<uint64_t>(wpf_bytes_, t->bytes)});
|
||||
}
|
||||
const char* k_comb() const { return ldb_ ? "moe_combine_sh2" : "moe_combine_sh"; }
|
||||
bool fuse_ = !(std::getenv("STRATA_METAL_FUSE") && std::string(std::getenv("STRATA_METAL_FUSE")) == "0");
|
||||
bool qdirty_ = false;
|
||||
// mac-kern's router_route (3f9ca996): the BF16 router GEMV and moe_route in one dispatch, beside the shared
|
||||
// expert's GEMVs (bitwise). Opt-in (STRATA_METAL_ROUTE=1): in this window, where the router GEMV already runs
|
||||
// beside the shared GEMVs, it measured no faster (T2 GPU 13.035 vs 13.027 ms; the fused kernel alone 0.615 ms vs
|
||||
// 0.566 for the two) and 0.3-0.5% slower end to end on the M5 Max.
|
||||
bool routef_ = std::getenv("STRATA_METAL_ROUTE") && std::string(std::getenv("STRATA_METAL_ROUTE")) == "1";
|
||||
mtl::Buffer rcounter_;
|
||||
// STRATA_METAL_KPROF=T (test tooling): once, after the 20th cached window of T tokens with the head, the GPU time
|
||||
// of each kernel class alone (the window's own dispatches of that pipeline + weight role, run by themselves)
|
||||
void kprof(const std::vector<mtl::Dispatch>& ds, int Tn);
|
||||
@@ -727,7 +705,6 @@ bool Window::init(const strata::core::ModelGeometry& g, strata::core::SessionSta
|
||||
if (const gguf::Tensor* e0 = f_.find("blk.0.ffn_gate_exps.weight")) gyro_dense_ = e0->type >= 144;
|
||||
lib_ = cx_->compile({"quant.metal", "units.metal", "moe.metal", "gemv.metal", "router.metal", "gr.metal", "gr2.metal",
|
||||
"glue.metal"}, metal_f16() ? "#define STRATA_METAL_F16 1\n" : "");
|
||||
rcounter_ = cx_->alloc(16);
|
||||
lgdn_ = cx_->compile({"gdn.metal"}, "#define NO_GGML_COMMON 1\n");
|
||||
lqsa_ = cx_->compile({"qsa.metal", "qsa2.metal"}, "#define NO_GGML_COMMON 1\n");
|
||||
lsmp_ = cx_->compile({"sample.metal"}, "#define NO_GGML_COMMON 1\n");
|
||||
@@ -1009,18 +986,6 @@ void Window::gr_read(std::vector<mtl::Dispatch>& ds, const std::string& pre, boo
|
||||
d3.grid[0] = N / 16;
|
||||
d3.tg[0] = 256;
|
||||
push(ds, d3);
|
||||
for (const auto& [off, bytes] : wpf_next_) {
|
||||
if (!wpf_flag_.handle) wpf_flag_ = cx_->alloc(4096);
|
||||
mtl::Dispatch p;
|
||||
p.p = P(lib_, "wpf", {}, "wpf");
|
||||
const uint32_t n16 = (uint32_t) (bytes / 16);
|
||||
p.args = {mtl::Arg::B(W_, off), mtl::Arg::B(wpf_flag_), mtl::Arg::V(n16)};
|
||||
p.tg[0] = 256;
|
||||
p.grid[0] = std::max<uint32_t>(1, std::min<uint32_t>(n16 / (256 * 4), 4096));
|
||||
nb_ = true; // beside gr_up
|
||||
push(ds, p);
|
||||
}
|
||||
wpf_next_.clear();
|
||||
cur_inj_ = 1 - cur_inj_;
|
||||
}
|
||||
|
||||
@@ -1056,14 +1021,13 @@ void Window::had_rows_ind(std::vector<mtl::Dispatch>& ds, int lib, const std::st
|
||||
void Window::layer(std::vector<mtl::Dispatch>& ds, int l, int ntok, int64_t pos0) {
|
||||
const std::string b = "blk." + std::to_string(l) + ".";
|
||||
// ---- the mixer: layer 1's read after the PLE is not applied again (its pending write was applied before it)
|
||||
if (l % 4 != 3) { wpf_add(b + "attn_qkv.weight"); wpf_add(b + "attn_gate.weight"); } else wpf_add(b + "attn_q.weight");
|
||||
gr_read(ds, b + "hc_attn", l > 0 && !(l == 1 && ple_on_), XA_, false, ntok);
|
||||
if (l % 4 != 3) {
|
||||
const int gi = l - l / 4;
|
||||
// a one-token window commits in the step; a longer one leaves the state and commit() folds the kept prefix
|
||||
// from the inputs kept per layer (qkvL_ ...)
|
||||
const bool self_commit = ntok == 1;
|
||||
const uint32_t cm = self_commit ? 1u : !partial_ ? 0u : win_head_ ? (0x100u | 1u) : 1u;
|
||||
const uint32_t cm = self_commit ? 1u : 0u;
|
||||
const size_t o_qkv = (size_t) gi * MAXT * 10240 * 4, o_ab = (size_t) gi * MAXT * 48 * 4, o_z = (size_t) gi * MAXT * 6144 * 4;
|
||||
gemv(ds, T(b + "attn_qkv.weight"), XA_, qkvL_, 10240, ntok, o_qkv);
|
||||
nb_ = true;
|
||||
@@ -1189,16 +1153,12 @@ void Window::layer(std::vector<mtl::Dispatch>& ds, int l, int ntok, int64_t pos0
|
||||
gemv(ds, T(b + "attn_output.weight"), XQ_, bo_, N, ntok);
|
||||
}
|
||||
// ---- the MoE
|
||||
wpf_add(b + "ffn_gate_inp.weight"); wpf_add(b + "ffn_gate_shexp.weight"); wpf_add(b + "ffn_up_shexp.weight");
|
||||
gr_read(ds, b + "hc_ffn", true, XF_, false, ntok);
|
||||
const gguf::Tensor& trt = T(b + "ffn_gate_inp.weight");
|
||||
const bool rf = routef_ && (trt.type == ref::BF16 || trt.type == ref::F32);
|
||||
if (!rf) {
|
||||
gemv(ds, trt, XF_, logits_, 512, ntok);
|
||||
nb_ = true;
|
||||
}
|
||||
gemv(ds, trt, XF_, logits_, 512, ntok);
|
||||
nb_ = true;
|
||||
if (stream_ && st_pred_ && l + 1 < 48) { // the next layer's routing guessed from this layer's FFN input
|
||||
nb_ = !rf; // beside the router GEMV; with router_route it is the first reader of XF_ here: a barrier
|
||||
nb_ = true; // beside the router GEMV
|
||||
gemv(ds, T("blk." + std::to_string(l + 1) + ".ffn_gate_inp.weight"), XF_, predL_, 512, ntok,
|
||||
(size_t) l * MAXT * 512 * 4);
|
||||
if (predL2_.ptr && l + 2 < 48) {
|
||||
@@ -1259,21 +1219,9 @@ void Window::layer(std::vector<mtl::Dispatch>& ds, int l, int ntok, int64_t pos0
|
||||
a.gx_gu = NFF / (gs * (32 / gl) * gr);
|
||||
a.gx_dn = N / (ds_ * (32 / dl) * dr);
|
||||
mtl::Dispatch r;
|
||||
if (rf) {
|
||||
// the router GEMV + moe_route (bitwise both), beside the shared expert's GEMVs
|
||||
r.p = P(lib_, "router_route_" + std::to_string(trt.type), mtl::Constants().i(0, ntok), "router_route_" + std::to_string(trt.type) + "#" + std::to_string(ntok));
|
||||
GemvArgs ga{(uint32_t) N, (uint32_t) nexp_, (uint32_t) ntok, (uint32_t) ref::row_bytes(trt.type, N), (uint32_t) N, 512};
|
||||
r.args = {mtl::Arg::B(W_, trt.offset), mtl::Arg::B(XF_.b, 0), mtl::Arg::B(XF_.b, XF_.plane_bytes() * 4),
|
||||
mtl::Arg::B(logits_), mtl::Arg::V(ga), mtl::Arg::B(ids_), mtl::Arg::B(wts_), mtl::Arg::B(plan_),
|
||||
mtl::Arg::V(a), mtl::Arg::B(rcounter_)};
|
||||
r.grid[0] = nexp_ / 16;
|
||||
r.tg[0] = 512;
|
||||
nb_ = true;
|
||||
} else {
|
||||
r.p = P(lib_, "moe_route", {}, "moe_route");
|
||||
r.args = {mtl::Arg::B(logits_), mtl::Arg::B(ids_), mtl::Arg::B(wts_), mtl::Arg::B(plan_), mtl::Arg::V(a)};
|
||||
r.tg[0] = 512;
|
||||
}
|
||||
r.p = P(lib_, "moe_route", {}, "moe_route");
|
||||
r.args = {mtl::Arg::B(logits_), mtl::Arg::B(ids_), mtl::Arg::B(wts_), mtl::Arg::B(plan_), mtl::Arg::V(a)};
|
||||
r.tg[0] = 512;
|
||||
push(ds, r);
|
||||
{
|
||||
GlueArgs ga{NFF, MAXT, N, TOPK};
|
||||
@@ -1283,7 +1231,7 @@ void Window::layer(std::vector<mtl::Dispatch>& ds, int l, int ntok, int64_t pos0
|
||||
d.grid[0] = NFF / 16;
|
||||
d.grid[1] = ntok;
|
||||
d.tg[0] = 16;
|
||||
nb_ = !rf; // with router_route the shared GEMVs are not behind a barrier yet: this one waits for them
|
||||
nb_ = true;
|
||||
push(ds, d);
|
||||
}
|
||||
const std::string gun = std::string(m16 ? "moe_gu16_" : "moe_gu_") + std::to_string(tg_.type) + "_" + std::to_string(gl) +
|
||||
@@ -1845,7 +1793,7 @@ void Window::stepprof(const std::vector<mtl::Dispatch>& ds, int Tn) {
|
||||
};
|
||||
if (c.rfind("moe_gu", 0) == 0) part = "experts gate+up";
|
||||
else if (c.rfind("moe_down", 0) == 0) part = "experts down";
|
||||
else if (c.rfind("moe_route", 0) == 0 || c.rfind("router_route", 0) == 0) part = "router";
|
||||
else if (c.rfind("moe_route", 0) == 0) part = "router";
|
||||
else if (c.rfind("moe_combine", 0) == 0) part = "combine";
|
||||
else if (c.rfind("gr_", 0) == 0) { part = "hc"; wb = c.rfind("gr_down", 0) == 0 ? tsize(0) + tsize(1) : c.rfind("gr_up", 0) == 0 ? tsize(0) : 0; }
|
||||
else if (c.rfind("gdn_", 0) == 0) part = "gdn other";
|
||||
@@ -1874,7 +1822,7 @@ void Window::stepprof(const std::vector<mtl::Dispatch>& ds, int Tn) {
|
||||
// the route dispatches (one per layer): the prefix that ends right after one holds that layer's routed ids
|
||||
std::vector<size_t> routes;
|
||||
for (size_t i = 0; i < n; ++i)
|
||||
if (pname.count(ds[i].p) && (pname[ds[i].p].rfind("moe_route", 0) == 0 || pname[ds[i].p].rfind("router_route", 0) == 0)) routes.push_back(i);
|
||||
if (pname.count(ds[i].p) && pname[ds[i].p].rfind("moe_route", 0) == 0) routes.push_back(i);
|
||||
const int reps = std::getenv("STRATA_METAL_STEPPROF_REPS") ? std::max(1, std::atoi(std::getenv("STRATA_METAL_STEPPROF_REPS"))) : 5;
|
||||
std::fprintf(stderr, "strata metal stepprof: T%d, %zu dispatches in %zu barrier groups, %zu layers routed, %d reps per prefix\n", Tn, n, m, routes.size(), reps);
|
||||
for (int i = 0; i < 40; ++i) cx_->run_graph(ds); // clocks up
|
||||
@@ -2017,7 +1965,7 @@ bool Window::finish(int Tn, const int32_t* tokens, int64_t pos0, bool head, int3
|
||||
}
|
||||
}
|
||||
logpos(Tn, tokens, pos0, head);
|
||||
if (Tn == 1 || partial_) hp_ = 1 - hp_; // the window committed its first row (or all of them) in the step
|
||||
if (Tn == 1) hp_ = 1 - hp_; // the window committed its first row in the step
|
||||
last_head_ = head;
|
||||
for (int t = 0; t < Tn; ++t) out[t] = head ? tokout_.as<int32_t>()[t] : -1;
|
||||
last_T_ = Tn;
|
||||
@@ -2038,7 +1986,7 @@ bool Window::finish(int Tn, const int32_t* tokens, int64_t pos0, bool head, int3
|
||||
}
|
||||
|
||||
void Window::fold(std::vector<mtl::Dispatch>& ds, int n_keep, int hp) {
|
||||
const size_t off = partial_ ? 1 : 0; // rows 1.. (row 0 was committed in the window's step)
|
||||
const size_t off = 0;
|
||||
for (int l = 0; l < 48; ++l) {
|
||||
if (l % 4 == 3) continue;
|
||||
const int gi = l - l / 4;
|
||||
@@ -2324,14 +2272,6 @@ bool Window::commit(int n_keep, std::string& err) {
|
||||
// when those positions run again, before anything reads them.
|
||||
static const bool lazy = !(std::getenv("STRATA_METAL_LAZY_FOLD") && std::string(std::getenv("STRATA_METAL_LAZY_FOLD")) == "0");
|
||||
int nf = n_keep;
|
||||
if (partial_) {
|
||||
if (!last_head_) { // the prompt's window committed every row in its step
|
||||
if (n_keep != last_T_) { err = "metal window: a prompt window commits all its rows"; return false; }
|
||||
nf = 0;
|
||||
} else {
|
||||
nf = n_keep - 1; // row 0 is in
|
||||
}
|
||||
}
|
||||
if (nf > 0) {
|
||||
if (lazy) {
|
||||
pend_fold_ = nf; // folded at the start of the next window (from the history parity before the turn)
|
||||
|
||||
@@ -21,7 +21,6 @@ The engine package folders are copied (never changed in place) before they are s
|
||||
from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import json
|
||||
import shutil
|
||||
import sys
|
||||
import tarfile
|
||||
|
||||
@@ -0,0 +1,5 @@
|
||||
# Packages the checks need on top of requirements.txt (tools/check, the CI workflows). Not part of the runtime or the app.
|
||||
# CMake and Ninja are in requirements-build.txt.
|
||||
pytest==9.1.1
|
||||
# the Terminal's screens are tested with Textual's test driver (tests/terminal); setup.py installs the same pin for the Terminal
|
||||
textual==8.2.8
|
||||
@@ -0,0 +1,26 @@
|
||||
# Layer-split reserve: the unused window-term proposal
|
||||
|
||||
Moved here from `include/strata/program/stage_reserve.hpp` and `tests/program/stage_reserve_test.cpp` in cleanup step 4.2.
|
||||
Nothing in the engine calls it; the default placement uses `reserve_mib` and `room_bytes` in the header, which are unchanged.
|
||||
|
||||
## What it is
|
||||
|
||||
Issue #3702: on an unequal pair (RTX 3060 + RTX 5060 Ti) the last card was left with 68 MiB after loading although 700 + 96 MiB were reserved for it, and the verify windows' graphs (captured on first use, 4-6 MiB each and once 42 MiB) then ran it out of memory. The proposal makes the window term of a later stage grow with the number of window sizes:
|
||||
|
||||
- `proposed_window_mib(max_window_t)` = `kWindowMib` + (sizes - 1) * `kGraphMib` (6 MiB) + `kInstantiateBurstMib` (48 MiB)
|
||||
- `proposed_reserve_mib(in, max_window_t)` = `reserve_mib(in)` with that term on a later stage
|
||||
|
||||
For `--spec 4` (windows of 1..6 tokens) the term is 174 MiB instead of 96. It changes cache capacity, so it can change speed and it is not a cleanup change.
|
||||
|
||||
## What it needs before it could become the default
|
||||
|
||||
The model gate `phaseA-tests/gh027/gates/3702-later-reserve.sh` (private lab folder, not in this repository), and an A/B on a mixed-card pair.
|
||||
|
||||
## Files
|
||||
|
||||
- `proposal.hpp`: the two functions and their two constants (text unchanged).
|
||||
- `proposal_test.cpp`: the checks that used to sit in `stage_reserve_test.cpp`.
|
||||
|
||||
Run the checks (from the repository root, any C++17 compiler):
|
||||
|
||||
c++ -std=c++17 -Iinclude -Iresearch/stage-reserve-proposal research/stage-reserve-proposal/proposal_test.cpp -o proposal_test && ./proposal_test
|
||||
@@ -0,0 +1,32 @@
|
||||
// research/stage-reserve-proposal/proposal.hpp - PROPOSAL, not used by the engine (moved out of
|
||||
// include/strata/program/stage_reserve.hpp in cleanup step 4.2; the text of both functions is unchanged).
|
||||
// See README.md in this folder.
|
||||
#pragma once
|
||||
|
||||
#include "strata/program/stage_reserve.hpp"
|
||||
|
||||
#include <algorithm>
|
||||
#include <cstdint>
|
||||
|
||||
namespace strata::program::stage_reserve {
|
||||
|
||||
constexpr int64_t kGraphMib = 6; ///< one window graph, captured on first use (log of #3702: 4, 6, 6)
|
||||
constexpr int64_t kInstantiateBurstMib = 48; ///< the largest single step seen (42 MiB in #3702): room kept for one more
|
||||
|
||||
/// PROPOSAL, unused: the window term of a later stage when the windows of every size 1..max_window_t can still be
|
||||
/// captured after loading: the first window's arena (kWindowMib already holds it), a graph for each further size, and
|
||||
/// one burst.
|
||||
inline int64_t proposed_window_mib(int max_window_t) {
|
||||
const int64_t sizes = std::max<int64_t>(max_window_t, 1);
|
||||
return kWindowMib + (sizes - 1) * kGraphMib + kInstantiateBurstMib;
|
||||
}
|
||||
|
||||
/// PROPOSAL, unused: reserve_mib with the proposed window term on a later stage.
|
||||
inline int64_t proposed_reserve_mib(const Inputs& in, int max_window_t) {
|
||||
Inputs p = in;
|
||||
int64_t r = reserve_mib(p);
|
||||
if (in.later) r += proposed_window_mib(max_window_t) - kWindowMib;
|
||||
return r;
|
||||
}
|
||||
|
||||
} // namespace strata::program::stage_reserve
|
||||
@@ -0,0 +1,36 @@
|
||||
// research/stage-reserve-proposal/proposal_test.cpp - the checks of the unused window-term proposal (moved out of
|
||||
// tests/program/stage_reserve_test.cpp in cleanup step 4.2; the assertions are unchanged). Not built by CMake.
|
||||
// Build and run by hand, see README.md in this folder.
|
||||
#include "proposal.hpp"
|
||||
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <string>
|
||||
|
||||
using namespace strata::program::stage_reserve;
|
||||
|
||||
namespace {
|
||||
int g_fail = 0;
|
||||
void check(bool ok, const std::string& what) {
|
||||
if (ok) return;
|
||||
std::fprintf(stderr, "FAIL: %s\n", what.c_str());
|
||||
++g_fail;
|
||||
}
|
||||
} // namespace
|
||||
|
||||
int main() {
|
||||
// the proposal: the window term grows with the number of window sizes (--spec 4 -> windows of 1..6 tokens)
|
||||
check(proposed_window_mib(1) == 96 + 48, "proposal: one window size = the arena and the burst");
|
||||
check(proposed_window_mib(6) == 96 + 5 * 6 + 48, "proposal: six sizes = 174 MiB");
|
||||
check(proposed_window_mib(0) == proposed_window_mib(1), "proposal: no negative sizes");
|
||||
{
|
||||
Inputs in;
|
||||
in.base_reserve_mib = 700;
|
||||
in.later = true;
|
||||
check(proposed_reserve_mib(in, 6) == 700 + 174, "proposed reserve on a later card, six sizes");
|
||||
in.later = false;
|
||||
check(proposed_reserve_mib(in, 6) == reserve_mib(in), "the first card is not changed by the proposal");
|
||||
}
|
||||
std::printf("proposal_test: %s\n", g_fail ? "FAIL" : "PASS");
|
||||
return g_fail ? 1 : 0;
|
||||
}
|
||||
@@ -10,7 +10,6 @@ used from the next start on.
|
||||
from __future__ import annotations
|
||||
|
||||
import json
|
||||
import os
|
||||
import shutil
|
||||
from pathlib import Path
|
||||
|
||||
|
||||
@@ -33,7 +33,6 @@ import base64
|
||||
import hashlib
|
||||
import hmac
|
||||
import codecs
|
||||
import ctypes
|
||||
import json
|
||||
import math
|
||||
import os
|
||||
|
||||
@@ -15,7 +15,6 @@ from __future__ import annotations
|
||||
import collections
|
||||
import ctypes
|
||||
import os
|
||||
import platform
|
||||
from strata_control import sysinfo
|
||||
import re
|
||||
import subprocess
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
// CUDA standalone of the sampler statistics harness: nvcc -include ds4_launch.h -I include -I src/dsv41 sampler_stat_main.cu
|
||||
// CUDA standalone of the sampler statistics harness. Build: cmake --build <dir> --target ds4_sampler_stat (optional, off by default; -DSTRATA_BUILD_DSV41=ON),
|
||||
// or by hand: nvcc -include ds4_launch.h -I include -I src/dsv41 sampler_stat_main.cu. Usage: ds4_sampler_stat ROWS.bin N_VOCAB N_SAMPLES OUTPREFIX (src/dsv41/sampler_stat.inc).
|
||||
#include "sampler_stat.inc"
|
||||
#include "vendor/sampler.cu"
|
||||
int main(int argc, char** argv) { return sampler_stat::run(argc, argv); }
|
||||
@@ -2181,7 +2181,7 @@ __global__ void __launch_bounds__(THREADS) gr_up_q8_kernel(GrMulti m, const floa
|
||||
|
||||
// STRATA_HC_Q8 on HIP gfx12: gr_up_q8_kernel with 32 columns / 512 threads a block, so a block owns a whole q8_1 group
|
||||
// (STRATA_QFUSE_HC: no counter, no fence; the same arithmetic per value as gr_up_q8_kernel)
|
||||
template <int T, int G = 1, bool PL = false>
|
||||
template <int T, int G = 1>
|
||||
__global__ void __launch_bounds__(Q32_THREADS) gr_up_q8_q32_kernel(GrMulti m, const float* __restrict__ part,
|
||||
const float* __restrict__ ssg) {
|
||||
__shared__ __align__(16) float lo[T][LR];
|
||||
@@ -2191,13 +2191,13 @@ __global__ void __launch_bounds__(Q32_THREADS) gr_up_q8_q32_kernel(GrMulti m, co
|
||||
const int d0 = blockIdx.x * Q32_COLS;
|
||||
constexpr int RPW8 = HC * Q32_COLS / Q32_WARPS; // 8 rows per warp
|
||||
constexpr int NG = RPW8 / G;
|
||||
constexpr int NB = PL ? 2 : 1;
|
||||
constexpr int NB = 1;
|
||||
float4 w[NB][G][3];
|
||||
float rvA[NB][G], wnA[NB][G], boA[NB][G], ipA[NB][G];
|
||||
bool applyA[NB][G];
|
||||
// G rows a warp are loaded together (weights, R, w_norm, bo, inj) before the first of them is computed and stored: the in-place
|
||||
// R_out store of one row otherwise ends the loads the compiler could issue for the next (G = 1 is the plain one-row loop). PL: the
|
||||
// next group is loaded before this one is computed, the first before the norm / partial-sum phases (their latency hides the weights').
|
||||
// R_out store of one row otherwise ends the loads the compiler could issue for the next (G = 1 is the plain one-row loop). (A
|
||||
// pipelined variant, STRATA_GR_Q32_PL, was removed in 0.2.13: no gain over G = 2.)
|
||||
// Every row's arithmetic is unchanged, so the outputs are bitwise the same.
|
||||
auto load_grp = [&](int qq0, int bf) {
|
||||
#pragma unroll
|
||||
@@ -2226,7 +2226,6 @@ __global__ void __launch_bounds__(Q32_THREADS) gr_up_q8_q32_kernel(GrMulti m, co
|
||||
}
|
||||
}
|
||||
};
|
||||
if constexpr (PL) load_grp(0, 0);
|
||||
if (t < T * HC) {
|
||||
const int k = t / HC, c = t - k * HC;
|
||||
float ss = 0.0f;
|
||||
@@ -2258,9 +2257,8 @@ __global__ void __launch_bounds__(Q32_THREADS) gr_up_q8_q32_kernel(GrMulti m, co
|
||||
__syncthreads();
|
||||
#pragma unroll
|
||||
for (int sg = 0; sg < NG; ++sg) {
|
||||
const int bf = PL ? (sg & 1) : 0;
|
||||
if constexpr (PL) { if (sg + 1 < NG) load_grp((sg + 1) * G, bf ^ 1); }
|
||||
else load_grp(sg * G, 0);
|
||||
const int bf = 0;
|
||||
load_grp(sg * G, 0);
|
||||
#pragma unroll
|
||||
for (int gg = 0; gg < G; ++gg) {
|
||||
const int r = warp + (sg * G + gg) * Q32_WARPS;
|
||||
@@ -2591,11 +2589,7 @@ template <int T> void launch_q8_t(const GrMulti& m, float* part, float* ssg, cud
|
||||
#if !STRATA_GR_FAST_BUILD
|
||||
if (m.a[0].q8_mixed != nullptr && qfuse_hc_local()) { // HIP (not gfx906): the block owns a q8_1 group
|
||||
static const int upg = [] { const char* v = std::getenv("STRATA_GR_Q32_G"); const int g = v ? std::atoi(v) : 1; return g == 2 || g == 4 || g == 8 ? g : 1; }();
|
||||
static const bool upl = [] { const char* v = std::getenv("STRATA_GR_Q32_PL"); return v != nullptr && std::atoi(v) != 0; }();
|
||||
if (upl && upg == 2) gr_up_q8_q32_kernel<T, 2, true><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else if (upl && upg == 4) gr_up_q8_q32_kernel<T, 4, true><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else if (upl) gr_up_q8_q32_kernel<T, 1, true><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else if (upg == 2) gr_up_q8_q32_kernel<T, 2><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
if (upg == 2) gr_up_q8_q32_kernel<T, 2><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else if (upg == 4) gr_up_q8_q32_kernel<T, 4><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else if (upg == 8) gr_up_q8_q32_kernel<T, 8><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
else gr_up_q8_q32_kernel<T, 1><<<Q32_BLOCKS, Q32_THREADS, 0, st>>>(m, part, ssg);
|
||||
|
||||
@@ -18018,8 +18018,6 @@ int main(int argc, char** argv) {
|
||||
std::vector<int32_t> drafts(DS, 0), window(WS), outv(WS);
|
||||
std::vector<float> dprob(DS, 0.0f);
|
||||
std::array<float, 8> dwin_s{}; // the probabilities the current window was chosen from (STRATA_SPEC_COST)
|
||||
strata::spec::AdaptiveGate df2_gate; // STRATA_DF2_GATE=1: this request's calibrated DFlash2 window (spec/window_rule.hpp)
|
||||
df2_gate.row_ms = strata::spec::df2_gate_row_ms();
|
||||
// STRATA_SPEC_TRACE=<file> in serve (bench only, off by default): per round "R p T a sfx ms k match | probs | drafts | lookup draft | committed tokens"
|
||||
static std::FILE* const df2_trace = [] { const char* tf = std::getenv("STRATA_SPEC_TRACE"); return tf != nullptr && tf[0] != '\0' ? std::fopen(tf, "a") : nullptr; }();
|
||||
static const bool trace_nopick = std::getenv("STRATA_SPEC_TRACE_NOPICK") != nullptr; // the lookup draft is logged but never taken
|
||||
@@ -18886,7 +18884,6 @@ int main(int argc, char** argv) {
|
||||
wr.chain_rows = chain_rows; wr.cap = S_mtp; wr.start = T; wr.spec_min_p = req_spec_min_p;
|
||||
wr.draft_conf = o.draft_conf; wr.dprob = dprob.data();
|
||||
T = strata::spec::choose_window(wr);
|
||||
if (req_df2 && draft_on && !first_window && strata::spec::df2_gate_on()) T = df2_gate.choose(dprob.data(), S_mtp);
|
||||
}
|
||||
if (first_window) T = merge_T > 0 ? merge_T : 1;
|
||||
// a repeat of earlier context (prompt lookup) where the MTP's own first guess agrees: the policy takes it
|
||||
@@ -19340,10 +19337,6 @@ int main(int argc, char** argv) {
|
||||
tr_k = 0;
|
||||
}
|
||||
if (spec_cost && use_mtp && timed_round && !from_sfx) policy.observe_probs(dwin_s.data(), T_mtp - 1, a);
|
||||
if (req_df2 && timed_round && strata::spec::df2_gate_on()) { // STRATA_DF2_GATE=1: what this round showed about the survival of its drafts, and what a token cost
|
||||
if (!from_sfx) df2_gate.observe(dwin_s.data(), T, a);
|
||||
df2_gate.observe_round(std::chrono::duration<double, std::milli>(Clock::now() - round0).count(), a + 1);
|
||||
}
|
||||
if (timed_round && !eos && chain_n == 0)
|
||||
policy.observe(from_sfx, T, a, sfx_match,
|
||||
std::chrono::duration<double, std::milli>(Clock::now() - round0).count());
|
||||
|
||||
@@ -103,21 +103,6 @@ int main() {
|
||||
check(!mtp_q_applies(true, false, false, false, false, 4), "greedy / mode off: no q");
|
||||
check(!mtp_q_applies(false, false, true, false, false, 4), "no MTP: no q");
|
||||
}
|
||||
{ // 4. the calibrated DFlash2 gate (STRATA_DF2_GATE=1)
|
||||
AdaptiveGate g;
|
||||
float sure[8] = {0.95f, 0.95f, 0.95f, 0.95f, 0.95f, 0.95f, 0.95f, 0.95f}, weak[8] = {0.2f, 0.9f, 0.9f, 0.9f, 0.9f, 0.9f, 0.9f, 0.9f};
|
||||
check(g.choose(sure, 6) == 6, "gate, no data, confident drafts: the whole window");
|
||||
check(g.choose(weak, 6) == 1, "gate, no data, a hopeless first draft: one row");
|
||||
for (int i = 0; i < 40; ++i) g.observe(sure, 6, 1); // drafts that keep failing at position 2 narrow the window for that bucket
|
||||
check(g.choose(sure, 6) <= 2, "gate learns that position 2 fails: the window narrows to 2 rows");
|
||||
AdaptiveGate g2;
|
||||
for (int i = 0; i < 40; ++i) g2.observe(sure, 6, 5);
|
||||
check(g2.choose(sure, 6) == 6, "gate, drafts always hold: the whole window");
|
||||
AdaptiveGate g3, g4;
|
||||
float mid[8] = {0.6f, 0.6f, 0.6f, 0.6f, 0.6f, 0.6f, 0.6f, 0.6f};
|
||||
g3.ms_per_tok = 4.0; g4.ms_per_tok = 14.0;
|
||||
check(g3.choose(mid, 6) <= g4.choose(mid, 6), "gate: a dearer token (14 ms) never gets a narrower window than a cheap one (4 ms)");
|
||||
}
|
||||
std::printf("%s\n", g_fail == 0 ? "window_rule_test: ok" : "window_rule_test: FAILED");
|
||||
return g_fail == 0 ? 0 : 1;
|
||||
}
|
||||
@@ -142,6 +142,24 @@ foreach(_n ${STRATA_SPLIT_CARDS})
|
||||
set_target_properties(${_t} PROPERTIES CUDA_RUNTIME_LIBRARY Shared)
|
||||
endforeach()
|
||||
|
||||
# the engine's kernel self-test on one card, with no model and no engine start (cleanup 4.4; tools/selftest_main.cpp). HIP only, off by default
|
||||
# (EXCLUDE_FROM_ALL: built by name). The engine runs the same src/selftest.cu before it loads a model on a card Strata has not run the split on;
|
||||
# this is the way to ask a card that question alone:
|
||||
# cmake --build <dir> --target strata-split-selftest && <dir>/src/split/strata-split-selftest [device] (exit 0: every test passed)
|
||||
if(_fnx_hip)
|
||||
list(GET STRATA_SPLIT_CARDS 0 _nst)
|
||||
set(_t strata-split-selftest)
|
||||
add_executable(${_t} EXCLUDE_FROM_ALL ${_sp}/tools/selftest_main.cpp ${_sp}/src/selftest.cu ${_sp}/src/moe_mma.cu ${_sp_kparts})
|
||||
set_source_files_properties(${_sp}/tools/selftest_main.cpp ${_sp}/src/selftest.cu ${_sp}/src/moe_mma.cu PROPERTIES LANGUAGE HIP)
|
||||
target_include_directories(${_t} PRIVATE ${_sp}/src/hip ${_sp}/src ${_sp}/tools ${STRATA_GGML_DIR}/ggml/src/ggml-cuda
|
||||
${STRATA_GGML_DIR}/ggml/src ${STRATA_GGML_DIR}/ggml/include)
|
||||
target_compile_definitions(${_t} PRIVATE FNX_NCARD=${_nst} FNX_ARCH=${STRATA_SPLIT_ARCH} FNX_27B=0 FNX_HOSTRING=1 FNX_HIP=1 GGML_USE_HIP=1
|
||||
FNX_DIR="${_sp}")
|
||||
target_compile_options(${_t} PRIVATE "$<$<COMPILE_LANGUAGE:HIP>:SHELL:-include ${_sp}/src/hip/fnx_hip.h>")
|
||||
target_link_libraries(${_t} PRIVATE strata_hip_runtime ${STRATA_HIP_BLAS_TARGETS})
|
||||
set_target_properties(${_t} PROPERTIES CXX_STANDARD 17 HIP_STANDARD 17 HIP_STANDARD_REQUIRED ON HIP_ARCHITECTURES "${_fnx_hip_archs}")
|
||||
endif()
|
||||
|
||||
# the CPU unit tests sFNX ships (split rule, ring plan, card checks, batch plan, rope, input parser): plain C++, no GPU
|
||||
if(STRATA_BUILD_TESTS)
|
||||
list(GET STRATA_SPLIT_CARDS 0 _n0)
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
// Standalone run of the engine's kernel self-test (src/selftest.cu) on one card: no model, no engine. Linked with the engine's kernel objects by
|
||||
// phaseA-tests/sfnx/hip/selftest_p2006.sh. Prints the card and the result; exit 0 when every test passed.
|
||||
// Standalone run of the engine's kernel self-test (src/selftest.cu) on one card: no model, no engine. Built by the optional CMake target strata-split-selftest
|
||||
// (src/split/CMakeLists.txt; HIP only, off by default). Prints the card and the result; exit 0 when every test passed.
|
||||
// selftest_main [device]
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
|
||||
@@ -288,7 +288,7 @@ def stream_chat(target: Target, spec: dict, model: str | None, *, timeout_s: flo
|
||||
raw["interrupted"] = True
|
||||
raw["error"] = {"kind": "cancelled", "message": "cancelled by the operator"}
|
||||
except (OSError, http.client.HTTPException, ValueError) as exc:
|
||||
if raw["timed_out"]:
|
||||
if raw["timed_out"] or isinstance(exc, TimeoutError): # the socket's own timeout can fire just before the timer does (Windows)
|
||||
raw["error"] = {"kind": "timeout", "message": f"no complete answer within {timeout_s:g} seconds"}
|
||||
else:
|
||||
raw["error"] = {"kind": "connection_" + type(exc).__name__, "message": str(exc)[:200] or type(exc).__name__}
|
||||
|
||||
@@ -39,7 +39,6 @@ class _VersionAction(argparse.Action):
|
||||
|
||||
def parser() -> argparse.ArgumentParser:
|
||||
ap = argparse.ArgumentParser(prog="strata", description="Strata serving and terminal clients")
|
||||
import strata_cli
|
||||
ap.add_argument("--version", action=_VersionAction, nargs=0,
|
||||
help="print the version and the interpreter this terminal runs on")
|
||||
sp = ap.add_subparsers(dest="command", required=True)
|
||||
|
||||
@@ -8,7 +8,6 @@ from __future__ import annotations
|
||||
|
||||
from typing import Any
|
||||
|
||||
from rich import box
|
||||
from rich.console import Group
|
||||
from rich.panel import Panel
|
||||
from rich.text import Text
|
||||
|
||||
@@ -19,7 +19,7 @@ import time
|
||||
from dataclasses import dataclass
|
||||
from pathlib import Path
|
||||
|
||||
from strata_cli.launch.settings import ACCESS_KINDS, LaunchSettings
|
||||
from strata_cli.launch.settings import LaunchSettings
|
||||
|
||||
TUNNEL_URL = re.compile(r"https://[a-z0-9-]+\.trycloudflare\.com")
|
||||
WINDOWS = os.name == "nt"
|
||||
|
||||
@@ -6,14 +6,12 @@ launch-aware extras: a Connect view (key 7) and a quit question (keep serving in
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import json
|
||||
import sys
|
||||
import threading
|
||||
from dataclasses import dataclass, field
|
||||
from pathlib import Path
|
||||
from typing import Any
|
||||
|
||||
from rich.text import Text
|
||||
from textual.app import App
|
||||
from textual.binding import Binding
|
||||
|
||||
@@ -431,7 +429,6 @@ def run_monitor(info: dict, *, keep_default: bool = True, color: str = "auto") -
|
||||
app = LaunchMonitor(endpoint, None, info={**info, "daily_check": True}, color=color) # (the quiet once-a-day check for a newer engine: this monitor only)
|
||||
app.run()
|
||||
if app.stop_requested:
|
||||
from strata_cli.launch.runner import Runner as R
|
||||
from strata_control import runtime
|
||||
try:
|
||||
runtime.stop(info["profile"], timeout=180)
|
||||
|
||||
@@ -7,9 +7,7 @@ import json
|
||||
import os
|
||||
import re
|
||||
import sys
|
||||
import threading
|
||||
from pathlib import Path
|
||||
from typing import Any
|
||||
|
||||
from strata_cli.launch import settings as S
|
||||
|
||||
@@ -592,7 +590,7 @@ def _printer(quiet: bool = False):
|
||||
|
||||
|
||||
def _stop_all(name: str | None) -> int:
|
||||
from strata_cli.launch import access, store
|
||||
from strata_cli.launch import access
|
||||
from strata_control import profiles, runtime
|
||||
targets = [name] if name else [p["id"] for p in profiles.list_profiles() if p["id"].startswith("launch-")]
|
||||
stopped = failed = 0
|
||||
@@ -1062,7 +1060,7 @@ def run_launch(a: argparse.Namespace) -> int:
|
||||
# ------------------------------------------------------------------------------------------------ connect
|
||||
|
||||
def run_connect(a: argparse.Namespace) -> int:
|
||||
from strata_cli.launch import access, harness
|
||||
from strata_cli.launch import harness
|
||||
machine = a.output == "json"
|
||||
if not a.harness:
|
||||
rows = [{"id": h.id, "title": h.title, "speaks": h.speaks, "status": h.status, "vendor_tested": h.tested or None} for h in harness.HARNESSES]
|
||||
|
||||
@@ -12,7 +12,6 @@ from __future__ import annotations
|
||||
|
||||
import json
|
||||
import os
|
||||
import re
|
||||
import shutil
|
||||
import time
|
||||
from dataclasses import dataclass, field
|
||||
|
||||
@@ -17,7 +17,6 @@ import textwrap
|
||||
import threading
|
||||
|
||||
from rich.text import Text
|
||||
from textual.binding import Binding
|
||||
from textual.widgets import Static
|
||||
|
||||
from strata_cli import nav
|
||||
|
||||
@@ -13,7 +13,6 @@ tests see the same thing. Nothing here prints or reads the keyboard. The steps r
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import json
|
||||
import os
|
||||
import re
|
||||
|
||||
@@ -8,11 +8,7 @@ of key hints. Esc goes back, Enter goes on, q leaves (and says what is kept). Ch
|
||||
from __future__ import annotations
|
||||
|
||||
import os
|
||||
import shutil
|
||||
import sys
|
||||
import threading
|
||||
from pathlib import Path
|
||||
from typing import Any
|
||||
|
||||
from rich.text import Text
|
||||
from textual.app import ComposeResult
|
||||
|
||||
@@ -10,7 +10,7 @@ from __future__ import annotations
|
||||
|
||||
import os
|
||||
import sys
|
||||
from dataclasses import dataclass, field
|
||||
from dataclasses import dataclass
|
||||
from typing import Any
|
||||
|
||||
from rich.text import Text
|
||||
|
||||
@@ -11,7 +11,6 @@ import shutil
|
||||
import sys
|
||||
import tempfile
|
||||
import threading
|
||||
import time
|
||||
from strata_control import atomic as _atomic_mod
|
||||
from pathlib import Path
|
||||
from typing import Any
|
||||
@@ -174,7 +173,7 @@ def _hw_hash(identity: dict[str, Any]) -> str:
|
||||
# --------------------------------------------------------------------------- tune auto
|
||||
|
||||
def auto(a: argparse.Namespace) -> int:
|
||||
from strata_control import autotune, engine_component, presets, tune_knobs, tune_machine, tuning
|
||||
from strata_control import autotune, engine_component, presets, tuning
|
||||
from strata_control.tune_session import WorkerSession
|
||||
|
||||
machine = a.output == "json"
|
||||
@@ -291,7 +290,7 @@ def _human_result(r: dict[str, Any], saved: dict[str, Any] | None) -> str:
|
||||
# --------------------------------------------------------------------------- presets
|
||||
|
||||
def presets_command(a: argparse.Namespace) -> int:
|
||||
from strata_control import autotune, presets, tune_knobs
|
||||
from strata_control import presets
|
||||
|
||||
machine = a.output == "json"
|
||||
cmd = a.presets_command
|
||||
|
||||
@@ -37,17 +37,15 @@ from __future__ import annotations
|
||||
|
||||
import hashlib
|
||||
import json
|
||||
import os
|
||||
import platform
|
||||
import re
|
||||
import shutil
|
||||
import stat
|
||||
import tarfile
|
||||
import time
|
||||
import zipfile
|
||||
from dataclasses import asdict, dataclass, field
|
||||
from dataclasses import asdict, dataclass
|
||||
from pathlib import Path, PurePosixPath
|
||||
from typing import Any, Iterable, Sequence
|
||||
from typing import Any, Iterable
|
||||
|
||||
SCHEMA_VERSION = 1
|
||||
|
||||
|
||||
@@ -176,7 +176,7 @@ def install_package(dest: Path, package: str | os.PathLike, *, emit: Callable[[s
|
||||
"""Install a DeepSeek engine package the person downloaded from the release page (a .tar.gz / .zip, or the unpacked folder) into `dest`.
|
||||
Checked before anything is installed: it is a DeepSeek engine package for this system, its manifest is signed by a key Strata trusts, and every
|
||||
file matches the hash in the manifest. Returns the new strata-dsv41, or None with a plain line through `emit`."""
|
||||
from strata_control import components, engine_component as ec, signing, updater
|
||||
from strata_control import engine_component as ec, signing, updater
|
||||
from strata_control.targets import host_platform
|
||||
|
||||
src, dest = Path(package).expanduser(), Path(dest)
|
||||
|
||||
@@ -20,7 +20,6 @@ setup.py, tools/strata_mcp.py and strata_cli/launch import it. Recommend, never
|
||||
from __future__ import annotations
|
||||
|
||||
import os
|
||||
import platform
|
||||
from strata_control import sysinfo
|
||||
import sys
|
||||
from typing import Any
|
||||
|
||||
@@ -12,7 +12,6 @@ measured figures and the person decides.
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
from typing import Any
|
||||
|
||||
MAX_CONVERSATIONS = 8 # the engine's limit (ds4_main.cpp: --parallel takes 1 to 8)
|
||||
|
||||
|
||||
@@ -23,7 +23,6 @@ from __future__ import annotations
|
||||
|
||||
import datetime as dt
|
||||
import json
|
||||
import re
|
||||
import threading
|
||||
from dataclasses import dataclass, field
|
||||
from pathlib import Path
|
||||
|
||||
@@ -35,9 +35,7 @@ would settle it. Recording the limit is the honest option; claiming enforcement
|
||||
from __future__ import annotations
|
||||
|
||||
import os
|
||||
import shutil
|
||||
from dataclasses import dataclass
|
||||
from pathlib import Path
|
||||
from typing import Any, Mapping
|
||||
from urllib.parse import urlparse
|
||||
|
||||
|
||||
@@ -495,7 +495,6 @@ def history_path() -> Path:
|
||||
|
||||
|
||||
def record_trials(trials: list[dict[str, Any]], *, hw_hash: str, model_digest: str | None, engine_sha: str | None) -> None:
|
||||
from strata_control import tune_knobs as k
|
||||
rows = []
|
||||
for t in trials:
|
||||
args = {}
|
||||
|
||||
@@ -33,7 +33,6 @@ from __future__ import annotations
|
||||
import dataclasses
|
||||
import pathlib
|
||||
import re
|
||||
from typing import Any
|
||||
|
||||
# What has been run with the host exchange, per make: {make: {architecture: card counts}}. NVIDIA: 2x RTX A6000 (boston, sm_86), 2x RTX PRO 6000 (srv118, sm_120).
|
||||
# AMD: 4x Radeon AI PRO R9700 (r9700, gfx1201; phaseA-tests/sfnx/RESULTS-HIP.md): 4 cards with UD-Q4_K_XL, 3 cards with the IQ3_S set, 2 cards with the IQ3_S set at 8K context and 2 slots only
|
||||
@@ -528,7 +527,6 @@ def record_launch_manifest(d: Decision, mode_: str, environ: dict | None = None)
|
||||
|
||||
|
||||
def _apply(cfg: dict, *, exe_for_cards, system, allow_untested, probe, ram, scan, data_dir, mode_: str = "auto") -> Decision:
|
||||
import os
|
||||
import pathlib
|
||||
import sys
|
||||
|
||||
|
||||
@@ -26,11 +26,9 @@ from __future__ import annotations
|
||||
import json
|
||||
import re
|
||||
import subprocess
|
||||
import sys
|
||||
import time
|
||||
import urllib.error
|
||||
import urllib.request
|
||||
from dataclasses import dataclass, field
|
||||
from pathlib import Path
|
||||
from typing import Any
|
||||
|
||||
|
||||
@@ -29,15 +29,13 @@ measurement, and neither is a run whose `produced` is zero.
|
||||
|
||||
from __future__ import annotations
|
||||
|
||||
import json
|
||||
import os
|
||||
import re
|
||||
import select
|
||||
import subprocess
|
||||
import time
|
||||
from dataclasses import dataclass, field
|
||||
from pathlib import Path
|
||||
from typing import Any, Iterable, Sequence
|
||||
from typing import Any, Sequence
|
||||
|
||||
#: A workload is fixed for the whole sweep: the same prompt, the same number of tokens to generate,
|
||||
#: the same sampling. Changing any of these between candidates invalidates the comparison.
|
||||
|
||||
@@ -27,13 +27,12 @@ from __future__ import annotations
|
||||
import hashlib
|
||||
import json
|
||||
import os
|
||||
import platform
|
||||
from strata_control import sysinfo
|
||||
import time
|
||||
import uuid
|
||||
from dataclasses import asdict, dataclass, field
|
||||
from pathlib import Path
|
||||
from typing import Any, Sequence
|
||||
from typing import Any
|
||||
|
||||
SCHEMA_VERSION = 1
|
||||
|
||||
|
||||
@@ -18,7 +18,7 @@ from __future__ import annotations
|
||||
|
||||
import sys
|
||||
from dataclasses import dataclass, field
|
||||
from typing import Any, Callable, Sequence
|
||||
from typing import Any, Sequence
|
||||
|
||||
from strata_control import user_args
|
||||
|
||||
|
||||
@@ -13,7 +13,6 @@ from __future__ import annotations
|
||||
import hashlib
|
||||
import json
|
||||
import os
|
||||
import platform
|
||||
from strata_control import sysinfo
|
||||
import re
|
||||
import shutil
|
||||
|
||||
+1
-1
@@ -4,7 +4,7 @@
|
||||
# taken from ../src unless this directory carries a migrated copy: `src/X.cu` -> `sycl/src/X.dp.cpp`,
|
||||
# `src/X.cpp` -> `sycl/src/X.cpp` (only host files that called the CUDA runtime were migrated). Headers resolve
|
||||
# the same way through the include order. Upstream files are never edited, so the port is the content of this
|
||||
# directory plus the tool that produced it (../../sycl/migrate.sh).
|
||||
# directory plus the tool that produced it (../../sycl/tools/migrate.sh).
|
||||
#
|
||||
# cmake -S sycl -B build-sycl -G Ninja -DCMAKE_CXX_COMPILER=icpx && cmake --build build-sycl
|
||||
#
|
||||
|
||||
File diff suppressed because it is too large.
Load diff
@@ -160,19 +160,6 @@ template <int SG> __dpct_inline__ float sg_sum(float x) { // SYCL port: warp_s
|
||||
x += sycl::permute_group_by_xor(sycl::ext::oneapi::this_work_item::get_sub_group(), x, offset);
|
||||
return x;
|
||||
}
|
||||
// SYCL port: the decode mmvq kernels as Xe2-native SIMD16 sub-groups (twice the registers per lane, same 128
|
||||
// threads per group) or as the CUDA-shaped 32-lane warp (default). STRATA_MMVQ_SG: 16 = all SIMD16, 1 = IQ4_XS
|
||||
// only, 2 = short outputs (n_out <= 1024) only, 3 = both. mmvq_sg_bench alone: SIMD16 wins on IQ4_XS (up to
|
||||
// 1.45x) and short outputs (up to 1.3x), loses 10-25% on the large K-quant projections.
|
||||
inline int mmvq_sg() {
|
||||
static const int v = std::getenv("STRATA_MMVQ_SG") ? std::atoi(std::getenv("STRATA_MMVQ_SG")) : 32;
|
||||
return v;
|
||||
}
|
||||
inline bool mmvq_sg16(bool iq4xs, int n_out) {
|
||||
const int v = mmvq_sg();
|
||||
return v == 16 || ((v & 1) && v < 4 && iq4xs) || ((v & 2) && v < 4 && n_out <= 1024);
|
||||
}
|
||||
|
||||
__dpct_inline__ float warp_max(float x) {
|
||||
#pragma unroll
|
||||
for (int offset = WARP / 2; offset > 0; offset >>= 1) {
|
||||
@@ -1632,87 +1619,7 @@ void native_mmvq_q6k_wide_kernel(const Q6KBlock* __restrict__ w, const Q81Block*
|
||||
if (lane == 0 && row0 + r < n_out) y[std::size_t(j) * n_out + row0 + r] = v;
|
||||
}
|
||||
}
|
||||
// SYCL port: the wide Q6_K kernel with its weight loads issued U steps ahead. The wide kernel's 2560-wide rows
|
||||
// are ~3 steps of load -> wait -> dot per sub-group, so each sub-group has one 16-byte load per lane in flight
|
||||
// at a time and the kernel is latency-bound (~146 GB/s). Here a sub-group first issues the ql/qh/scale loads
|
||||
// of U steps (4U blocks), then does their dots; the activation reads stay inline (they are L1 hits, shared by
|
||||
// every row). One row per sub-group, WARPS rows per group.
|
||||
template <int NCOLS, int U, bool A2 = false>
|
||||
void native_mmvq_q6k_pf_kernel(const Q6KBlock* __restrict__ w, const Q81Block* __restrict__ x,
|
||||
float* __restrict__ y, int n_in, int n_out) {
|
||||
auto item = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
|
||||
const int lane = int(item.get_local_id(2)), warp = int(item.get_local_id(1));
|
||||
const int row = int(item.get_group(2)) * WARPS + warp;
|
||||
if (row >= n_out) return;
|
||||
const int blocks_per_row = n_in / QK, x_stride = n_in / Q8K;
|
||||
const int g = lane & 7, sub = lane >> 3;
|
||||
const int vh_shift = 2 * ((g & 3) >> 1);
|
||||
const int qh4_idx = 2 * (g >> 2) + (g & 1);
|
||||
const int scale_offset = 8 * (g >> 2) + (g & 3);
|
||||
const int bq8_offset = 4 * (g >> 2) + ((g & 3) >> 1);
|
||||
const int u_int4 = g & 1;
|
||||
const Q6KBlock* wr = w + std::size_t(row) * blocks_per_row;
|
||||
float acc[NCOLS] = {};
|
||||
for (int kb0 = 0; kb0 < blocks_per_row; kb0 += 4 * U) {
|
||||
sycl::int4 ql4[U], qh4[U]; float dsc0[U], dsc1[U]; bool ok[U];
|
||||
#pragma unroll
|
||||
for (int u = 0; u < U; ++u) { // all U steps' weight loads first
|
||||
const int kbx = kb0 + 4 * u + sub;
|
||||
ok[u] = kbx < blocks_per_row;
|
||||
const Q6KBlock* b = wr + (ok[u] ? kbx : 0);
|
||||
if constexpr (A2) { ql4[u] = load16_a2(b->ql + 16 * g); qh4[u] = load16_a2(b->qh + 16 * qh4_idx); }
|
||||
else { ql4[u] = reinterpret_cast<const sycl::int4*>(b->ql)[g];
|
||||
qh4[u] = reinterpret_cast<const sycl::int4*>(b->qh)[qh4_idx]; }
|
||||
const float d = b->d;
|
||||
dsc0[u] = d * (float) b->scales[scale_offset]; dsc1[u] = d * (float) b->scales[scale_offset + 4];
|
||||
}
|
||||
#pragma unroll
|
||||
for (int u = 0; u < U; ++u) {
|
||||
if (!ok[u]) continue;
|
||||
const int kby = (kb0 + 4 * u + sub) * (QK / Q8K);
|
||||
const int vl[4] = {ql4[u].x(), ql4[u].y(), ql4[u].z(), ql4[u].w()};
|
||||
const int vh[4] = {qh4[u].x() >> vh_shift, qh4[u].y() >> vh_shift, qh4[u].z() >> vh_shift, qh4[u].w() >> vh_shift};
|
||||
int vi0[4], vi1[4];
|
||||
#pragma unroll
|
||||
for (int p = 0; p < 4; ++p) {
|
||||
vi0[p] = q6_sub32((vl[p] & 0x0f0f0f0f) | ((vh[p] << 4) & 0x30303030));
|
||||
vi1[p] = q6_sub32(((vl[p] >> 4) & 0x0f0f0f0f) | (((vh[p] >> 4) << 4) & 0x30303030));
|
||||
}
|
||||
#pragma unroll
|
||||
for (int j = 0; j < NCOLS; ++j) {
|
||||
const Q81Block* xb = x + std::size_t(j) * x_stride + kby + bq8_offset;
|
||||
sycl::int4 a0, a1;
|
||||
if constexpr (A2) { // Q8_1 qs at offset 4 of a 36-byte block: four int loads (as the wide kernel)
|
||||
const int* p0 = reinterpret_cast<const int*>(xb[0].qs) + 4 * u_int4;
|
||||
const int* p1 = reinterpret_cast<const int*>(xb[2].qs) + 4 * u_int4;
|
||||
a0 = sycl::int4(p0[0], p0[1], p0[2], p0[3]);
|
||||
a1 = sycl::int4(p1[0], p1[1], p1[2], p1[3]);
|
||||
} else {
|
||||
a0 = reinterpret_cast<const sycl::int4*>(xb[0].qs)[u_int4];
|
||||
a1 = reinterpret_cast<const sycl::int4*>(xb[2].qs)[u_int4];
|
||||
}
|
||||
int s0 = 0, s1 = 0;
|
||||
s0 = strata::dp4a(vi0[0], a0.x(), s0); s0 = strata::dp4a(vi0[1], a0.y(), s0);
|
||||
s0 = strata::dp4a(vi0[2], a0.z(), s0); s0 = strata::dp4a(vi0[3], a0.w(), s0);
|
||||
s1 = strata::dp4a(vi1[0], a1.x(), s1); s1 = strata::dp4a(vi1[1], a1.y(), s1);
|
||||
s1 = strata::dp4a(vi1[2], a1.z(), s1); s1 = strata::dp4a(vi1[3], a1.w(), s1);
|
||||
acc[j] += (float) xb[0].ds[0] * dsc0[u] * (float) s0 + (float) xb[2].ds[0] * dsc1[u] * (float) s1;
|
||||
}
|
||||
}
|
||||
}
|
||||
auto sg = item.get_sub_group();
|
||||
#pragma unroll
|
||||
for (int j = 0; j < NCOLS; ++j) {
|
||||
float v = acc[j];
|
||||
#pragma unroll
|
||||
for (int o = WARP / 2; o > 0; o >>= 1) v += sycl::permute_group_by_xor(sg, v, o);
|
||||
if (lane == 0) y[std::size_t(j) * n_out + row] = v;
|
||||
}
|
||||
}
|
||||
inline int q6k_pf() { // STRATA_MMVQ_PF: weight-load steps issued ahead in the wide Q6_K kernel (0 = off)
|
||||
static const int v = std::getenv("STRATA_MMVQ_PF") ? std::atoi(std::getenv("STRATA_MMVQ_PF")) : 0;
|
||||
return v;
|
||||
}
|
||||
// (STRATA_MMVQ_PF, a prefetching wide Q6_K kernel, and STRATA_MMVQ_SG, SIMD16 sub-group variants of the decode mmvq, were removed in 0.2.13: no gain in the engine.)
|
||||
inline int q6k_wide_rpw() {
|
||||
static const int v = std::getenv("STRATA_MMVQ_WIDE_RPW") ? std::atoi(std::getenv("STRATA_MMVQ_WIDE_RPW")) : 1;
|
||||
return v;
|
||||
@@ -1735,22 +1642,6 @@ void launch_q6k_wide(const void* weights, const void* x_q8_1, float* y, int n_in
|
||||
const auto* w = static_cast<const Q6KBlock*>(weights);
|
||||
const auto* x = static_cast<const Q81Block*>(x_q8_1);
|
||||
const int rpw = q6k_wide_rpw();
|
||||
if (const int pf = q6k_pf(); pf > 0) {
|
||||
const unsigned blocks = unsigned((std::size_t(n_out) + WARPS - 1) / WARPS);
|
||||
const sycl::range<3> t(1, WARPS, WARP);
|
||||
#define STRATA_Q6PF(UU) if (pf == UU && q6k_a2()) { \
|
||||
strata::kprof::kp(s, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class native_mmvq_q6k_pf_a2, dpct_kernel_scalar<NCOLS>, dpct_kernel_scalar<UU>>>( \
|
||||
sycl::nd_range<3>(sycl::range(1, 1, blocks) * t, t), \
|
||||
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] { native_mmvq_q6k_pf_kernel<NCOLS, UU, true>(w, x, y, n_in, n_out); }); \
|
||||
return; } \
|
||||
if (pf == UU) { \
|
||||
strata::kprof::kp(s, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class native_mmvq_q6k_pf, dpct_kernel_scalar<NCOLS>, dpct_kernel_scalar<UU>>>( \
|
||||
sycl::nd_range<3>(sycl::range(1, 1, blocks) * t, t), \
|
||||
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(32)]] { native_mmvq_q6k_pf_kernel<NCOLS, UU>(w, x, y, n_in, n_out); }); \
|
||||
return; }
|
||||
STRATA_Q6PF(2) STRATA_Q6PF(3) STRATA_Q6PF(4) STRATA_Q6PF(6)
|
||||
#undef STRATA_Q6PF
|
||||
}
|
||||
if (rpw == 1 && q6k_a2()) {
|
||||
const unsigned blocks = unsigned((std::size_t(n_out) + WARPS - 1) / WARPS);
|
||||
if constexpr (NCOLS > 1) {
|
||||
@@ -1786,14 +1677,6 @@ void launch_q6k_wide(const void* weights, const void* x_q8_1, float* y, int n_in
|
||||
});
|
||||
return;
|
||||
}
|
||||
if (rpw == 1 && mmvq_sg16(false, n_out)) {
|
||||
constexpr int NWG = WARPS * WARP / 16;
|
||||
const unsigned blocks = unsigned((std::size_t(n_out) + NWG - 1) / NWG);
|
||||
strata::kprof::kp(s, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class native_mmvq_q6k_wide16, dpct_kernel_scalar<NCOLS>>>(
|
||||
sycl::nd_range<3>(sycl::range(1, 1, blocks) * sycl::range(1, NWG, 16), sycl::range(1, NWG, 16)),
|
||||
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(16)]] { native_mmvq_q6k_wide_kernel<NCOLS, 1, 16>(w, x, y, n_in, n_out); });
|
||||
return;
|
||||
}
|
||||
// RP rows per sub-group share each column's activation loads (the rows' own lanes, k order and butterfly: every
|
||||
// output the same bits); with the aligned loads (A2) unless STRATA_MMVQ_A2=0
|
||||
#define STRATA_Q6W(RP) if (rpw == RP) { \
|
||||
@@ -2209,25 +2092,6 @@ void launch_multi_n(const void *weights, const void *x_q8_1, float *y, int n_in,
|
||||
});
|
||||
return;
|
||||
}
|
||||
if (mmvq_sg16(std::is_same_v<F, IQ4XSTraits>, n_out)) { // same 128 threads as 8 SIMD16 sub-groups
|
||||
constexpr int NW16 = WARPS * WARP / 16;
|
||||
const sycl::range<3> t16(1, NW16, 16);
|
||||
if (n_in / F::DIV < F::BPI) {
|
||||
const unsigned blocks = unsigned((std::size_t(n_out) + WARPS - 1) / WARPS);
|
||||
strata::kprof::kp(s, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class native_mmvq_multi16_small, F, dpct_kernel_scalar<NCOLS>>>(
|
||||
sycl::nd_range<3>(sycl::range(1, 1, blocks) * t16, t16),
|
||||
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(16)]] {
|
||||
native_mmvq_multi_kernel<F, NCOLS, NW16, WARPS, false, false, 16>(w, x, y, n_in, n_out);
|
||||
});
|
||||
} else {
|
||||
strata::kprof::kp(s, __FILE__, __func__, __LINE__)->template parallel_for<dpct_kernel_name<class native_mmvq_multi16, F, dpct_kernel_scalar<NCOLS>>>(
|
||||
sycl::nd_range<3>(sycl::range(1, 1, unsigned(n_out)) * t16, t16),
|
||||
[=](sycl::nd_item<3>) [[sycl::reqd_sub_group_size(16)]] {
|
||||
native_mmvq_multi_kernel<F, NCOLS, NW16, 1, false, false, 16>(w, x, y, n_in, n_out);
|
||||
});
|
||||
}
|
||||
return;
|
||||
}
|
||||
const dpct::dim3 threads(WARP, WARPS);
|
||||
// #783 PR-h (stuchapin909): two rows per block for every K, not only small K - each row keeps its own partial sums and
|
||||
// reduction, so a row's result does not depend on its neighbour; STRATA_NO_MMVQ_ROWS2=1 keeps one row for large K
|
||||
|
||||
@@ -40,44 +40,7 @@ struct TileA { uint32_t x[4]; };
|
||||
struct TileB { uint32_t x[2]; };
|
||||
struct TileC { float x[4]={0.0f,0.0f,0.0f,0.0f}; };
|
||||
#if !defined(__HIPCC__)
|
||||
#if 0 // SYCL: inline PTX (mma/ldmatrix/cp.async) - the XMX port is pending; see tools/fixups.py
|
||||
__dpct_inline__ void load_a(TileA &a, const float *p) {
|
||||
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
|
||||
const float *src = p + (item_ct1.get_local_id(2) % 16) * STRIDE +
|
||||
(item_ct1.get_local_id(2) / 16) * 4;
|
||||
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
|
||||
asm volatile("ldmatrix.sync.aligned.m8n8.x4.b16 {%0,%1,%2,%3}, [%4];"
|
||||
: "=r"(a.x[0]), "=r"(a.x[1]), "=r"(a.x[2]), "=r"(a.x[3])
|
||||
: "l"(src));
|
||||
#else
|
||||
dpct::experimental::matrix::ldmatrix((uintptr_t)src, &a.x[0], &a.x[1],
|
||||
&a.x[2], &a.x[3]);
|
||||
#endif
|
||||
}
|
||||
__dpct_inline__ void load_b(TileB &b, const float *p) {
|
||||
auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
|
||||
const float *src = p + (item_ct1.get_local_id(2) % 8) * STRIDE +
|
||||
((item_ct1.get_local_id(2) / 8) * 4) % 8;
|
||||
#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
|
||||
asm volatile("ldmatrix.sync.aligned.m8n8.x2.b16 {%0,%1}, [%2];"
|
||||
: "=r"(b.x[0]), "=r"(b.x[1])
|
||||
: "l"(src));
|
||||
#else
|
||||
dpct::experimental::matrix::ldmatrix((uintptr_t)src, &b.x[0], &b.x[1]);
|
||||
#endif
|
||||
}
|
||||
__dpct_inline__ void mma(TileC &c, const TileA &a, const TileB &b) {
|
||||
// Deliberately no cvt.rn.tf32: pinned mma.cuh passes raw F32 bits directly.
|
||||
/*
|
||||
DPCT1053: Migration of device assembly code is not supported.
|
||||
*/
|
||||
asm("mma.sync.aligned.m16n8k8.row.col.f32.tf32.tf32.f32 {%0,%1,%2,%3}, "
|
||||
"{%4,%5,%6,%7}, {%8,%9}, {%0,%1,%2,%3};"
|
||||
: "+f"(c.x[0]), "+f"(c.x[1]), "+f"(c.x[2]), "+f"(c.x[3])
|
||||
: "r"(a.x[0]), "r"(a.x[1]), "r"(a.x[2]), "r"(a.x[3]), "r"(b.x[0]),
|
||||
"r"(b.x[1]));
|
||||
}
|
||||
#endif
|
||||
// (no load_a/load_b/mma here: the inline-PTX tensor-core path is src/kernels/cuda/native_qsa_score.cu, and the SYCL port scores with the scalar path below; sycl/tools/fixups.py keeps this block out)
|
||||
/*
|
||||
DPCT1110: The total declared local variable size in device function
|
||||
score_kernel exceeds 128 bytes and may cause high register pressure. Consult
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
// sycl/src/kernels/mmvq_sg_bench.cpp - SYCL port: the decode mmvq kernels on the Coder's dense shapes and types.
|
||||
// mmvq_sg_bench (run with STRATA_MMVQ_SG=16 and =32; prints us per call and a checksum of the outputs)
|
||||
// mmvq_sg_bench (prints us per call and a checksum of the outputs; the SIMD16 variants it once compared, STRATA_MMVQ_SG, were removed in 0.2.13)
|
||||
#include "strata/kernels/native_mmvq.hpp"
|
||||
#include <sycl/sycl.hpp>
|
||||
#include <dpct/dpct.hpp>
|
||||
|
||||
@@ -4538,34 +4538,7 @@ int main(int argc, char **argv) try {
|
||||
strata::emulated_cc()
|
||||
? " (STRATA_EMULATE_CC: a test mode, the card is emulated)"
|
||||
: "");
|
||||
#if 0 // the SYCL port has no CUDA runtime to mismatch (#542): dpct::get_major_version is the device version
|
||||
#if !defined(STRATA_HIP_GFX906) // a CUDA toolkit mismatch; the gfx906 HIP build
|
||||
// has no DPCT_COMPAT_RT_VERSION
|
||||
{ // #542: a build whose libcudart is older than its headers (a CUDA 13 kit with a dangling libcudart.so that
|
||||
// CMake resolved to the system's CUDA 12 one) reads cudaDeviceProp shifted - silently, and slowly
|
||||
int rt = 0;
|
||||
/*
|
||||
DPCT1043: The version-related API is different in SYCL. An
|
||||
initial code was generated, but you need to adjust it.
|
||||
*/
|
||||
if (DPCT_CHECK_ERROR(rt = dpct::get_major_version(
|
||||
dpct::get_current_device())) == 0 &&
|
||||
rt / 1000 != DPCT_COMPAT_RT_VERSION / 1000)
|
||||
std::fprintf(
|
||||
stderr,
|
||||
"strata generate: WARNING: this engine was compiled with "
|
||||
"CUDA %d.%d headers but "
|
||||
"loaded a CUDA %d.%d runtime (libcudart): GPU properties "
|
||||
"can read wrong and some "
|
||||
"kernels go unused. Rebuild it against one toolkit (cmake "
|
||||
"-DCUDAToolkit_ROOT=<the "
|
||||
"toolkit>, with its libcudart.so present) (#542)\n",
|
||||
DPCT_COMPAT_RT_VERSION / 1000,
|
||||
DPCT_COMPAT_RT_VERSION % 1000 / 10, rt / 1000,
|
||||
rt % 1000 / 10);
|
||||
}
|
||||
#endif
|
||||
#endif // #if 0
|
||||
// (the CUDA runtime/headers mismatch warning, #542, lives in src/program/generate.cpp; a SYCL port has no CUDA runtime to mismatch; sycl/tools/fixups.py keeps this block out)
|
||||
#endif
|
||||
const std::string e = strata::core::device_code_error();
|
||||
if (!e.empty()) {
|
||||
|
||||
+30
-2
@@ -1,9 +1,9 @@
|
||||
#!/usr/bin/env python3
|
||||
"""Hand fixes on top of the SYCLomatic output (sycl/migrate.sh). Idempotent: re-run after a re-migration.
|
||||
"""Hand fixes on top of the SYCLomatic output (sycl/tools/migrate.sh). Idempotent: re-run after a re-migration.
|
||||
|
||||
Each entry is a thing dpct 2025.3 got wrong or could not do, with the reason. Upstream files are never touched.
|
||||
"""
|
||||
import re, sys, pathlib
|
||||
import re, pathlib
|
||||
root = pathlib.Path(__file__).resolve().parents[1]
|
||||
changed = 0
|
||||
|
||||
@@ -408,3 +408,31 @@ def version_flag(s):
|
||||
" }\n"
|
||||
"#if (defined(STRATA_USE_HIP)", 1)
|
||||
edit("src/program/generate.cpp", version_flag)
|
||||
|
||||
# 14. No permanently excluded blocks in the port (cleanup 4.1). Three `#if 0` regions used to sit in the migrated sources: the AMD-tuned S26/S27
|
||||
# expert kernels in iq_kernels.dp.cpp (997 lines), the inline-PTX tile helpers that fixup 11 switches off in native_qsa_score.dp.cpp, and
|
||||
# the #542 CUDA-runtime mismatch warning in generate.cpp. Each is dead text here and has a canonical implementation in src/ (the .cu
|
||||
# file or src/program/generate.cpp). A re-migration or a 3-way merge could bring them back, so this removes each region (depth-aware:
|
||||
# through its matching #endif) and leaves one comment. Idempotent: the marker line is gone after the first run.
|
||||
# qsa_prompt_attn.dp.cpp and qsa_select.dp.cpp keep their `#elif 0` arms: those sit in a live #if chain and are not part of this item.
|
||||
def strip_dead_region(marker, comment, next_line=None):
|
||||
def fn(s):
|
||||
L = s.split("\n")
|
||||
for i, line in enumerate(L):
|
||||
if line.lstrip().startswith(marker) and (next_line is None or next_line in L[i + 1]):
|
||||
depth = 0
|
||||
for j in range(i, len(L)):
|
||||
t = L[j].lstrip()
|
||||
if re.match(r"#\s*(if|ifdef|ifndef)\b", t): depth += 1
|
||||
elif re.match(r"#\s*endif\b", t):
|
||||
depth -= 1
|
||||
if depth == 0:
|
||||
return "\n".join(L[:i] + [comment] + L[j + 1:])
|
||||
elif depth == 1 and re.match(r"#\s*(else|elif)\b", t):
|
||||
return s # an #else arm would be live code: leave it for a person to read
|
||||
return s
|
||||
return s
|
||||
return fn
|
||||
edit("src/kernels/cuda/iq_kernels.dp.cpp", strip_dead_region("#if 0", "// (the S26/S27 kernels and their launchers are not here: the canonical code is src/kernels/cuda/iq_kernels.cu, STRATA_EXPERT_V2 / STRATA_EXPERT_V2K; sycl/tools/fixups.py keeps this block out)", "S26 STRATA_EXPERT_V2=1"))
|
||||
edit("src/kernels/cuda/native_qsa_score.dp.cpp", strip_dead_region("#if 0 // SYCL: inline PTX", "// (no load_a/load_b/mma here: the inline-PTX tensor-core path is src/kernels/cuda/native_qsa_score.cu, and the SYCL port scores with the scalar path below; sycl/tools/fixups.py keeps this block out)"))
|
||||
edit("src/program/generate.cpp", strip_dead_region("#if 0 // the SYCL port has no CUDA runtime", "// (the CUDA runtime/headers mismatch warning, #542, lives in src/program/generate.cpp; a SYCL port has no CUDA runtime to mismatch; sycl/tools/fixups.py keeps this block out)"))
|
||||
@@ -1,6 +1,6 @@
|
||||
#!/usr/bin/env bash
|
||||
# CUDA -> SYCL migration of the Strata engine with SYCLomatic (dpct), run inside strata-sycl-dev.
|
||||
# sycl/migrate.sh [repo] [out] repo defaults to the Strata_B70 checkout beside this dir
|
||||
# sycl/tools/migrate.sh [repo] [out] repo defaults to the Strata_B70 checkout beside this dir
|
||||
# Writes a compilation database first (there is no nvcc here, so CMake cannot make one), then migrates every
|
||||
# CUDA-touching source into <out> mirroring the tree. Files dpct leaves alone are not copied - the port's
|
||||
# CMake falls back to the original for those, so the diff against upstream stays the port itself.
|
||||
|
||||
@@ -7,8 +7,11 @@ import unittest
|
||||
from pathlib import Path
|
||||
|
||||
ROOT = Path(__file__).resolve().parents[2]
|
||||
BASH = shutil.which("bash") # the full path: on Windows a bare "bash" is found in System32 first, which is the WSL launcher
|
||||
SCRIPT = ROOT / "apps" / "desktop" / "scripts" / "make-tester-readme.sh"
|
||||
TEMPLATES = ROOT / "apps" / "desktop" / "tester"
|
||||
# the README names the app version, which make-tester-readme.sh reads from Cargo.toml
|
||||
APP_VERSION = next(line.split('"')[1] for line in (ROOT / "apps" / "desktop" / "src-tauri" / "Cargo.toml").read_text().splitlines() if line.startswith("version = "))
|
||||
NOTICE = "Strata opened in your browser because the app window couldn't start on this computer."
|
||||
|
||||
|
||||
@@ -16,11 +19,16 @@ def sha(path: Path) -> str:
|
||||
return hashlib.sha256(path.read_bytes()).hexdigest()
|
||||
|
||||
|
||||
def arg(value):
|
||||
# bash on Windows (Git for Windows, MSYS2, Cygwin) reads C:/dir/file, not C:\dir\file: the backslashes would be eaten
|
||||
return value.as_posix() if isinstance(value, Path) else str(value)
|
||||
|
||||
|
||||
def run(*args, check=True):
|
||||
return subprocess.run(["bash", str(SCRIPT), *map(str, args)], capture_output=True, text=True, check=check)
|
||||
return subprocess.run([BASH, arg(SCRIPT), *map(arg, args)], capture_output=True, text=True, check=check)
|
||||
|
||||
|
||||
@unittest.skipIf(shutil.which("bash") is None or shutil.which("sha256sum") is None, "needs bash and sha256sum")
|
||||
@unittest.skipIf(BASH is None or shutil.which("sha256sum") is None, "needs bash and sha256sum")
|
||||
class TesterReadmeTests(unittest.TestCase):
|
||||
def setUp(self):
|
||||
self.tmp = Path(tempfile.mkdtemp(prefix="strata-readme-"))
|
||||
@@ -39,7 +47,7 @@ class TesterReadmeTests(unittest.TestCase):
|
||||
text = self.out.read_text()
|
||||
self.assertIn(sha(app), text)
|
||||
self.assertNotIn("@APPIMAGE@", text)
|
||||
for needle in ("libfuse2t64", "Ubuntu 26.04", "--browser", "--quit", "--save-logs", "Strata 0.2.3", NOTICE):
|
||||
for needle in ("libfuse2t64", "Ubuntu 26.04", "--browser", "--quit", "--save-logs", f"Strata {APP_VERSION}", NOTICE):
|
||||
self.assertIn(needle, text)
|
||||
run("--verify", self.out, "--files-dir", self.files)
|
||||
|
||||
|
||||
@@ -1,6 +1,6 @@
|
||||
// tests/program/stage_reserve_test.cpp - #3702: the later-stage reserve arithmetic of the layer split search.
|
||||
// reserve_mib / room_bytes must equal generate.cpp's `stage_room` expression as it was written there; the proposal
|
||||
// functions are pinned too (they are not used by the default placement).
|
||||
// reserve_mib / room_bytes must equal generate.cpp's `stage_room` expression as it was written there. (The unused
|
||||
// window-term proposal and its checks are in research/stage-reserve-proposal/.)
|
||||
#include "strata/program/stage_reserve.hpp"
|
||||
|
||||
#include <cstdint>
|
||||
@@ -72,18 +72,6 @@ int main() {
|
||||
in.later = true;
|
||||
check(reserve_mib(in) == 1496, "the workaround reserve: 1400 + 96");
|
||||
}
|
||||
// the proposal: the window term grows with the number of window sizes (--spec 4 -> windows of 1..6 tokens)
|
||||
check(proposed_window_mib(1) == 96 + 48, "proposal: one window size = the arena and the burst");
|
||||
check(proposed_window_mib(6) == 96 + 5 * 6 + 48, "proposal: six sizes = 174 MiB");
|
||||
check(proposed_window_mib(0) == proposed_window_mib(1), "proposal: no negative sizes");
|
||||
{
|
||||
Inputs in;
|
||||
in.base_reserve_mib = 700;
|
||||
in.later = true;
|
||||
check(proposed_reserve_mib(in, 6) == 700 + 174, "proposed reserve on a later card, six sizes");
|
||||
in.later = false;
|
||||
check(proposed_reserve_mib(in, 6) == reserve_mib(in), "the first card is not changed by the proposal");
|
||||
}
|
||||
std::printf("stage_reserve_test: %ld cases, %s\n", cases, g_fail ? "FAIL" : "PASS");
|
||||
return g_fail ? 1 : 0;
|
||||
}
|
||||
@@ -969,6 +969,9 @@ class EngineProcessTests(unittest.TestCase):
|
||||
|
||||
def wait(self, timeout=None):
|
||||
return 0
|
||||
|
||||
def kill(self): # stop_engine's last step on Windows, when taskkill (mocked below) left the process running
|
||||
pass
|
||||
seen = {}
|
||||
|
||||
def popen(command, **kwargs):
|
||||
|
||||
@@ -152,7 +152,7 @@ class DraftLayerInTheDataDir(Base):
|
||||
|
||||
def test_plan_still_asks_for_the_layer_when_no_data_dir_has_it(self):
|
||||
folder = self.model_files(self.base / "ext" / "models" / "iq3_xxs")
|
||||
with mock.patch.object(Catalog, "data_dir", return_value=self.base / "emptydata"):
|
||||
with mock.patch.object(Catalog, "data_dir", return_value=self.base / "emptydata"), mock.patch.object(Runner, "intel_chosen", lambda self: False): # an Intel Arc in the box that runs the test would plan its own engine
|
||||
out = self.plan_text("--model-path", str(folder), "--data-dir", str(self.base / "data"), "--vision", "off")
|
||||
self.assertIn("About 5.2 GB to download: the draft layer (5.2 GB).", out)
|
||||
|
||||
@@ -323,7 +323,9 @@ class PlanOnBusyCards(Base):
|
||||
r = Runner(s, catalog=catalog({**HW_12, "ram_gb": 125.0, "engine_vendor": "nvidia",
|
||||
"gpus": [{"index": i, "name": "NVIDIA RTX PRO 6000", "vram_gb": 96.0, "usable": True, "vendor": "nvidia"} for i in range(2)]}))
|
||||
for p in (mock.patch.object(SE, "probe_cards", lambda idx=None: probe), mock.patch.object(SE, "ram_gib", lambda: (500.0, 400.0)),
|
||||
mock.patch("strata_control.drafting.detect_hardware", lambda: ("cuda", ["sm_120"]))):
|
||||
mock.patch("strata_control.drafting.detect_hardware", lambda: ("cuda", ["sm_120"])),
|
||||
mock.patch("sys.platform", "linux"), # the all-cards split rule asks the host OS (Linux only): this fake PC is Linux
|
||||
mock.patch.object(Runner, "intel_chosen", lambda self: False)): # and NVIDIA: the real box's Intel Arc must not make the launch an Intel one
|
||||
p.start()
|
||||
self.addCleanup(p.stop)
|
||||
return r
|
||||
@@ -368,7 +370,7 @@ class AppPlanWhileItsEngineRuns(Base):
|
||||
data = self.base / "d"
|
||||
data.mkdir(exist_ok=True)
|
||||
busy = [card(0, free=8192), card(1, free=8192)]
|
||||
with mock.patch.object(SE, "probe_cards", lambda idx=None: busy), mock.patch.object(SE, "ram_gib", lambda: (500.0, 400.0)), mock.patch("strata_control.drafting.detect_hardware", lambda: ("cuda", ["sm_120"])):
|
||||
with mock.patch.object(SE, "probe_cards", lambda idx=None: busy), mock.patch.object(SE, "ram_gib", lambda: (500.0, 400.0)), mock.patch("strata_control.drafting.detect_hardware", lambda: ("cuda", ["sm_120"])), mock.patch("sys.platform", "linux"), mock.patch.object(Runner, "intel_chosen", lambda self: False): # the split rule asks the host OS (Linux only) and the make: this fake PC is Linux and NVIDIA
|
||||
code, out, err = run_cli(["drafting", "plan", "--family", "qwen", "--data-dir", str(data), "--mode", "auto", "--output", "json", *extra])
|
||||
self.assertEqual(code, 0, err)
|
||||
return json.loads(out)
|
||||
|
||||
@@ -830,6 +830,10 @@ class AccessRules(unittest.TestCase):
|
||||
"https://calm-river-1234.trycloudflare.com")
|
||||
|
||||
|
||||
# a person at Git Bash: no CI variable (a CI runner sets one, and strata launch then never asks) and no STRATA_NO_PROMPT
|
||||
GIT_BASH_ENV = {"MSYSTEM": "MINGW64", "CI": "", "STRATA_NO_PROMPT": ""}
|
||||
|
||||
|
||||
class FlagsCommand(unittest.TestCase):
|
||||
"""Every `strata launch` flag, run through the real parser with the computer faked and nothing installed."""
|
||||
|
||||
@@ -992,7 +996,7 @@ class FlagsCommand(unittest.TestCase):
|
||||
|
||||
def test_git_bash_without_a_tty_goes_on_in_plain_questions(self):
|
||||
fake = Result("launch-x", ["http://127.0.0.1:8080"], None, None, [], "/x")
|
||||
with mock.patch.dict(os.environ, {"MSYSTEM": "MINGW64"}), mock.patch("builtins.input", return_value="y"), mock.patch.object(Runner, "run", return_value=fake) as run:
|
||||
with mock.patch.dict(os.environ, GIT_BASH_ENV), mock.patch("builtins.input", return_value="y"), mock.patch.object(Runner, "run", return_value=fake) as run:
|
||||
code, out, err = run_cli(["launch"])
|
||||
self.assertEqual(code, 0, err)
|
||||
run.assert_called_once()
|
||||
@@ -1000,13 +1004,13 @@ class FlagsCommand(unittest.TestCase):
|
||||
self.assertIn("winpty strata launch", out + err)
|
||||
|
||||
def test_git_bash_answering_no_changes_nothing(self):
|
||||
with mock.patch.dict(os.environ, {"MSYSTEM": "MINGW64"}), mock.patch("builtins.input", return_value="n"), mock.patch.object(Runner, "run", side_effect=AssertionError("must not run")):
|
||||
with mock.patch.dict(os.environ, GIT_BASH_ENV), mock.patch("builtins.input", return_value="n"), mock.patch.object(Runner, "run", side_effect=AssertionError("must not run")):
|
||||
code, out, err = run_cli(["launch"])
|
||||
self.assertEqual(code, 0, err)
|
||||
self.assertIn("Nothing was changed", out + err)
|
||||
|
||||
def test_a_pipe_with_nobody_to_answer_still_says_what_to_name(self):
|
||||
with mock.patch.dict(os.environ, {"MSYSTEM": "MINGW64"}), mock.patch("builtins.input", side_effect=EOFError), mock.patch.object(Runner, "run", side_effect=AssertionError("must not run")):
|
||||
with mock.patch.dict(os.environ, GIT_BASH_ENV), mock.patch("builtins.input", side_effect=EOFError), mock.patch.object(Runner, "run", side_effect=AssertionError("must not run")):
|
||||
code, out, err = run_cli(["launch"])
|
||||
self.assertEqual(code, 2)
|
||||
self.assertIn("--preset NAME", err)
|
||||
|
||||
@@ -28,6 +28,10 @@ class LaunchDeepSeek(unittest.TestCase):
|
||||
p = mock.patch.object(Catalog, "hardware", lambda s, fresh=False: dict(self.hw))
|
||||
p.start()
|
||||
self.addCleanup(p.stop)
|
||||
# several AMD cards share a DeepSeek model on Linux only (dsv41_fit.run_cards asks the host OS); these four R9700 are a Linux PC
|
||||
p = mock.patch("strata_control.dsv41_fit._on_windows", lambda: False)
|
||||
p.start()
|
||||
self.addCleanup(p.stop)
|
||||
|
||||
def plan(self, *args):
|
||||
code, out, err = run_cli(["launch", "--plan", "--output", "json", *args])
|
||||
|
||||
@@ -220,7 +220,8 @@ class MonitorBundleTests(unittest.TestCase):
|
||||
self.assertEqual(target.stat().st_mode & 0o077, 0)
|
||||
self.assertEqual((target / bundle.name).stat().st_mode & 0o077, 0)
|
||||
portable = subprocess.run([sys.executable, "-I", str(target / bundle.name), "logs", "--help"],
|
||||
cwd=directory, capture_output=True, text=True, check=True)
|
||||
cwd=directory, capture_output=True, text=True, check=True,
|
||||
encoding="utf-8", errors="replace") # (the help text has non-ASCII characters; Windows would read it as cp1252)
|
||||
self.assertIn("--follow", portable.stdout)
|
||||
extra = target / "user-notes.txt"
|
||||
extra.write_text("keep me")
|
||||
|
||||
@@ -31,7 +31,7 @@ def tree(root: Path) -> dict:
|
||||
for p in sorted(root.rglob("*")):
|
||||
if p.is_file():
|
||||
st = p.stat()
|
||||
out[str(p.relative_to(root))] = (st.st_size, st.st_mtime_ns)
|
||||
out[p.relative_to(root).as_posix()] = (st.st_size, st.st_mtime_ns)
|
||||
return out
|
||||
|
||||
|
||||
|
||||
@@ -10,6 +10,7 @@ import sys
|
||||
import tempfile
|
||||
import unittest
|
||||
from pathlib import Path
|
||||
from unittest import mock
|
||||
|
||||
sys.path.insert(0, str(Path(__file__).resolve().parents[2]))
|
||||
|
||||
@@ -120,6 +121,10 @@ class ElevatedRunLockoutTests(unittest.TestCase):
|
||||
self.tmp = tempfile.TemporaryDirectory()
|
||||
self.addCleanup(self.tmp.cleanup)
|
||||
self.root = Path(self.tmp.name)
|
||||
# repair_access only touches the user's own folders (user_roots): the scratch folder is the "profile" here, wherever TEMP points
|
||||
env = mock.patch.dict(os.environ, {"STRATA_HOME": str(self.root)})
|
||||
env.start()
|
||||
self.addCleanup(env.stop)
|
||||
|
||||
@staticmethod
|
||||
def lock_out(path):
|
||||
|
||||
@@ -23,6 +23,7 @@ from strata_cli.launch import settings as S
|
||||
from strata_cli.launch.runner import Runner
|
||||
from strata_control import profiles, split_engine as SE
|
||||
from tests.terminal.launch_support import HW_12, Home, catalog
|
||||
import serve.server # noqa: F401 imported before any test patches sys.platform: serve/lifeline.py runs Linux-only libc code when it is imported as Linux
|
||||
|
||||
A6000 = {"name": "NVIDIA RTX A6000", "vram_gb": 48.0, "usable": True, "vendor": "nvidia"}
|
||||
HW_2XA6000 = {**HW_12, "ram_gb": 125.0, "disk_free_gb": 900.0, "engine_vendor": "nvidia", "gpus": [{**A6000, "index": i} for i in range(2)]}
|
||||
@@ -51,6 +52,14 @@ class Base(unittest.TestCase):
|
||||
def setUp(self):
|
||||
self.home = Home().__enter__()
|
||||
self.addCleanup(self.home.__exit__, None, None, None)
|
||||
# the all-cards split rule asks the host OS and is Linux only (strata_control/split_engine.expected): these fake cards are a Linux PC
|
||||
plat = mock.patch("sys.platform", "linux")
|
||||
plat.start()
|
||||
self.addCleanup(plat.stop)
|
||||
# and NVIDIA: an Intel Arc in the box that runs the test would make the launch an Intel one (Runner.intel_chosen reads the real machine)
|
||||
intel = mock.patch.object(Runner, "intel_chosen", lambda self: False)
|
||||
intel.start()
|
||||
self.addCleanup(intel.stop)
|
||||
|
||||
def profile_for(self, cfg: dict, *, parallel=None, hw=None, four_card=None) -> dict:
|
||||
"""What Runner.ensure_profile saves as the launch profile's overrides for this written run config."""
|
||||
|
||||
@@ -350,7 +350,9 @@ class Monitor(unittest.TestCase):
|
||||
@contextlib.contextmanager
|
||||
def tempfile_home():
|
||||
import tempfile
|
||||
with tempfile.TemporaryDirectory(prefix="strata-notice-") as d:
|
||||
# A check the app started in the background may still hold .updates/update.lock when the case ends (a file in use cannot be deleted on
|
||||
# Windows); the folder is scratch, so a leftover is not a failure of the case.
|
||||
with tempfile.TemporaryDirectory(prefix="strata-notice-", ignore_cleanup_errors=True) as d:
|
||||
yield d
|
||||
|
||||
|
||||
|
||||
+1
-1
@@ -6,7 +6,7 @@
|
||||
arm = name=engine[|env;env[|extra args]]. Summary: median per case of the LAST repeat of each load (the first read of a
|
||||
prompt in an engine is low), ratio to the first arm, and whether the greedy tokens equal the first arm's (same round).
|
||||
"""
|
||||
import argparse, json, os, statistics, subprocess, sys
|
||||
import argparse, json, statistics, subprocess, sys
|
||||
|
||||
W = "/mnt/models/strata-test/arc037"
|
||||
|
||||
|
||||
@@ -26,7 +26,6 @@ import collections
|
||||
import mmap
|
||||
import os
|
||||
import pathlib
|
||||
import struct
|
||||
import sys
|
||||
import time
|
||||
|
||||
@@ -39,7 +38,7 @@ sys.path.insert(0, HERE)
|
||||
sys.path.insert(0, os.path.join(REPO, "ref"))
|
||||
|
||||
import gguf_reader as G # noqa: E402
|
||||
from gguf_writer import dequantize_q2_0 # noqa: E402
|
||||
from gguf_writer import dequantize_q2_0 # noqa: E402,F401 (kept: the packaged Control's bundle list follows this import)
|
||||
from load import _dequant_flat # noqa: E402
|
||||
|
||||
# default: $STRATA_SHARD1, else the development layout (<Strata>/Q2_0 beside <Strata>/Public/Engine)
|
||||
|
||||
+504
@@ -0,0 +1,504 @@
|
||||
#!/usr/bin/env python3
|
||||
"""One command for the fast checks: python tools/check (Python 3.8+, standard library only)
|
||||
|
||||
Runs, each in its own step with its own timeout, in a scratch folder (build trees, temp files, logs):
|
||||
|
||||
cmake-cpu CMake preset `cpu-check` (CMakePresets.json): configure, build, CTest. No GPU, no model, no network.
|
||||
py-tests pytest tests/ (without tests/snapshot, which has its own step)
|
||||
py-serve pytest serve/
|
||||
py-tools pytest tools/test_*.py
|
||||
py-optional tools/dflash2, tools/dsv41, tools/trellis tests; need torch and are skipped without it
|
||||
host-typecheck apps/host: npm run typecheck
|
||||
host-test apps/host: npm test (a per-test timeout, so a hang fails one test instead of the run)
|
||||
host-build apps/host: npm run build
|
||||
snapshot pytest tests/snapshot (the config snapshot of the four entry points; needs Node 24)
|
||||
|
||||
Prints one table: passed / failed / skipped / not run, per step, with the reason for every skip and every "not run".
|
||||
A step that cannot run because a tool is missing (Node, cmake, a compiler, pytest, torch) is "not run" and says which tool;
|
||||
it is never reported as a pass. The exit code is 1 when anything failed or timed out, and 0 otherwise; --strict makes a
|
||||
"not run" step exit 2 too. Nothing here needs a GPU, a model or the real Strata install.
|
||||
|
||||
python tools/check everything above
|
||||
python tools/check --only cmake-cpu,py-tools
|
||||
python tools/check --skip host-build --json result.json
|
||||
python tools/check --list
|
||||
Backend, model and browser tests are separate commands (see the CI workflows in .github/workflows); they are not part of this one.
|
||||
"""
|
||||
import argparse
|
||||
import glob
|
||||
import json
|
||||
import os
|
||||
import re
|
||||
import shutil
|
||||
import signal
|
||||
import subprocess
|
||||
import sys
|
||||
import tempfile
|
||||
import time
|
||||
import xml.etree.ElementTree as ET
|
||||
|
||||
ROOT = os.path.dirname(os.path.dirname(os.path.abspath(__file__)))
|
||||
WIN = os.name == "nt"
|
||||
HOST = os.path.join(ROOT, "apps", "host")
|
||||
NODE_MIN = (24, 12)
|
||||
|
||||
# step name -> default timeout in seconds (explicit, per step; --timeout-scale multiplies them)
|
||||
TIMEOUTS = {"cmake-cpu": 1800, "py-tests": 3600, "py-serve": 1800, "py-tools": 1200, "py-optional": 900,
|
||||
"host-typecheck": 600, "host-test": 1800, "host-build": 600, "snapshot": 900}
|
||||
STEPS = list(TIMEOUTS)
|
||||
|
||||
|
||||
class Result:
|
||||
def __init__(self, name):
|
||||
self.name, self.status, self.reason = name, "not run", ""
|
||||
self.passed = self.failed = self.skipped = 0
|
||||
self.seconds = 0.0
|
||||
self.failures = [] # test ids or the failing command
|
||||
self.skips = {} # reason -> count
|
||||
self.log = ""
|
||||
|
||||
def row(self):
|
||||
return {"step": self.name, "status": self.status, "passed": self.passed, "failed": self.failed,
|
||||
"skipped": self.skipped, "seconds": round(self.seconds, 1), "reason": self.reason,
|
||||
"failures": self.failures[:200], "skips": self.skips, "log": self.log}
|
||||
|
||||
|
||||
class Scratch:
|
||||
def __init__(self, base, keep):
|
||||
self.keep = keep
|
||||
self.created = base is None
|
||||
self.dir = os.path.abspath(base) if base else tempfile.mkdtemp(prefix="strata-check-")
|
||||
os.makedirs(os.path.join(self.dir, "tmp"), exist_ok=True)
|
||||
|
||||
def path(self, *p):
|
||||
return os.path.join(self.dir, *p)
|
||||
|
||||
def env(self, extra=None):
|
||||
env = dict(os.environ)
|
||||
tmp = self.path("tmp")
|
||||
env.update({"TMP": tmp, "TEMP": tmp, "TMPDIR": tmp, "PYTHONUTF8": "1", "PYTHONDONTWRITEBYTECODE": "1",
|
||||
"PYTHONUNBUFFERED": "1"}) # CI is left as the caller has it: some code asks whether a person can answer (strata_cli/launch/cmd.py)
|
||||
env.update(extra or {})
|
||||
return env
|
||||
|
||||
def cleanup(self):
|
||||
# only a folder this run created is removed
|
||||
if self.created and not self.keep:
|
||||
shutil.rmtree(self.dir, ignore_errors=True)
|
||||
|
||||
|
||||
def run_cmd(cmd, cwd, env, timeout, logfile, shell=False):
|
||||
"""Run one command to a log file. Returns (exit code or None on timeout, seconds). On timeout the whole tree of
|
||||
the child is killed (only this child's own tree)."""
|
||||
t0 = time.time()
|
||||
kw = {}
|
||||
if WIN:
|
||||
kw["creationflags"] = 0x00000200 # CREATE_NEW_PROCESS_GROUP
|
||||
else:
|
||||
kw["start_new_session"] = True
|
||||
with open(logfile, "ab") as log:
|
||||
log.write(("\n$ %s\n" % (cmd if isinstance(cmd, str) else " ".join(cmd))).encode())
|
||||
log.flush()
|
||||
try:
|
||||
p = subprocess.Popen(cmd, cwd=cwd, env=env, stdout=log, stderr=subprocess.STDOUT, shell=shell, **kw)
|
||||
except OSError as e:
|
||||
log.write(("could not start: %s\n" % e).encode())
|
||||
return -1, 0.0
|
||||
try:
|
||||
rc = p.wait(timeout=timeout)
|
||||
except subprocess.TimeoutExpired:
|
||||
if WIN:
|
||||
subprocess.call(["taskkill", "/PID", str(p.pid), "/T", "/F"], stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL)
|
||||
else:
|
||||
try:
|
||||
os.killpg(p.pid, signal.SIGTERM)
|
||||
time.sleep(3)
|
||||
os.killpg(p.pid, signal.SIGKILL)
|
||||
except (ProcessLookupError, PermissionError):
|
||||
pass
|
||||
p.wait()
|
||||
log.write(("\n[tools/check: timed out after %d s, killed]\n" % timeout).encode())
|
||||
return None, time.time() - t0
|
||||
return rc, time.time() - t0
|
||||
|
||||
|
||||
def tail(logfile, n=25):
|
||||
try:
|
||||
with open(logfile, "rb") as f:
|
||||
return f.read().decode("utf-8", "replace").splitlines()[-n:]
|
||||
except OSError:
|
||||
return []
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------------------------------------- tools
|
||||
|
||||
def node_version():
|
||||
node = shutil.which("node")
|
||||
if not node:
|
||||
return None, "Node.js is not installed (needs >= %d.%d)" % NODE_MIN
|
||||
try:
|
||||
out = subprocess.run([node, "--version"], capture_output=True, text=True, timeout=30).stdout.strip()
|
||||
v = tuple(int(x) for x in re.match(r"v(\d+)\.(\d+)", out).groups())
|
||||
except Exception as e: # noqa: BLE001
|
||||
return None, "node --version failed (%s)" % e
|
||||
if v < NODE_MIN:
|
||||
return None, "Node.js %s is too old (needs >= %d.%d)" % (out, NODE_MIN[0], NODE_MIN[1])
|
||||
return node, ""
|
||||
|
||||
|
||||
def have_module(name):
|
||||
r = subprocess.run([sys.executable, "-c", "import importlib.util,sys;sys.exit(0 if importlib.util.find_spec(%r) else 1)" % name],
|
||||
capture_output=True)
|
||||
return r.returncode == 0
|
||||
|
||||
|
||||
_vc_env = {}
|
||||
|
||||
|
||||
def msvc_env(base_env):
|
||||
"""On Windows the compiler is found through Visual Studio's vcvars64.bat when cl is not already on PATH."""
|
||||
if shutil.which("cl", path=base_env.get("PATH")):
|
||||
return base_env, ""
|
||||
if "env" in _vc_env:
|
||||
return _vc_env["env"], _vc_env["why"]
|
||||
cands = []
|
||||
pf = [os.environ.get("ProgramFiles(x86)", r"C:\Program Files (x86)"), os.environ.get("ProgramFiles", r"C:\Program Files")]
|
||||
vswhere = os.path.join(pf[0], "Microsoft Visual Studio", "Installer", "vswhere.exe")
|
||||
if os.path.exists(vswhere):
|
||||
out = subprocess.run([vswhere, "-latest", "-products", "*", "-requires", "Microsoft.VisualStudio.Component.VC.Tools.x86.x64",
|
||||
"-property", "installationPath"], capture_output=True, text=True).stdout.strip()
|
||||
if out:
|
||||
cands.append(os.path.join(out.splitlines()[0], "VC", "Auxiliary", "Build", "vcvars64.bat"))
|
||||
for root in pf:
|
||||
for edition in ("BuildTools", "Community", "Professional", "Enterprise"):
|
||||
cands.append(os.path.join(root, "Microsoft Visual Studio", "2022", edition, "VC", "Auxiliary", "Build", "vcvars64.bat"))
|
||||
for bat in cands:
|
||||
if os.path.exists(bat):
|
||||
out = subprocess.run('cmd /c ""%s" >nul && set"' % bat, capture_output=True, text=True, shell=True).stdout
|
||||
env = dict(base_env)
|
||||
for line in out.splitlines():
|
||||
if "=" in line and not line.startswith("="):
|
||||
k, v = line.split("=", 1)
|
||||
env[k] = v
|
||||
_vc_env["env"], _vc_env["why"] = env, ""
|
||||
return env, ""
|
||||
_vc_env["env"], _vc_env["why"] = base_env, "no C++ compiler: cl is not on PATH and Visual Studio 2022's vcvars64.bat was not found"
|
||||
return base_env, _vc_env["why"]
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------------------------------------- parsing
|
||||
|
||||
def parse_junit(path, res, prefix=""):
|
||||
try:
|
||||
tree = ET.parse(path)
|
||||
except (OSError, ET.ParseError):
|
||||
return False
|
||||
for case in tree.iter("testcase"):
|
||||
name = (case.get("classname") or "") + "::" + (case.get("name") or "")
|
||||
kids = {c.tag for c in case}
|
||||
if "skipped" in kids:
|
||||
sk = case.find("skipped")
|
||||
msg = (sk.get("message") or "").strip()
|
||||
if msg.lower().startswith("collection skipped") and (sk.text or "").strip(): # a module-level skip: the reason is in the text
|
||||
msg = (sk.text or "").strip().splitlines()[-1].replace("Skipped: ", "").strip()
|
||||
reason = (msg or (sk.text or "")).strip().splitlines()
|
||||
reason = reason[0][:140] if reason else "no reason given"
|
||||
if reason.lower() in ("true", "skipped", "skip"):
|
||||
reason = "(skipped in the test file without a reason: " + (case.get("name") or "?")[:80] + ")"
|
||||
res.skipped += 1
|
||||
res.skips[reason] = res.skips.get(reason, 0) + 1
|
||||
elif "failure" in kids or "error" in kids:
|
||||
res.failed += 1
|
||||
res.failures.append(prefix + name)
|
||||
else:
|
||||
res.passed += 1
|
||||
return True
|
||||
|
||||
|
||||
def finish(res, rc, timed_out, log):
|
||||
res.log = log
|
||||
if timed_out:
|
||||
res.status = "failed"
|
||||
res.reason = "timed out after %d s" % round(res.seconds)
|
||||
res.failures.append("timeout")
|
||||
elif rc != 0 or res.failed:
|
||||
res.status = "failed"
|
||||
if not res.failures:
|
||||
res.failures.append("exit code %s" % rc)
|
||||
res.reason = " | ".join(l.strip() for l in tail(log, 4))[:300]
|
||||
else:
|
||||
res.status = "passed"
|
||||
return res
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------------------------------------- steps
|
||||
|
||||
def step_cmake(ctx):
|
||||
r = Result("cmake-cpu")
|
||||
cmake, ninja = shutil.which("cmake"), shutil.which("ninja")
|
||||
if not cmake:
|
||||
r.reason = "cmake is not installed"
|
||||
return r
|
||||
if not ninja:
|
||||
r.reason = "ninja is not installed (the cpu-check preset uses the Ninja generator)"
|
||||
return r
|
||||
env = ctx.scratch.env()
|
||||
if WIN:
|
||||
env, why = msvc_env(env)
|
||||
if why:
|
||||
r.reason = why
|
||||
return r
|
||||
elif not any(shutil.which(c) for c in ("c++", "g++", "clang++")):
|
||||
r.reason = "no C++ compiler (c++, g++ or clang++) on PATH"
|
||||
return r
|
||||
build = ctx.scratch.path("build-cpu-check")
|
||||
log = ctx.scratch.path("cmake-cpu.log")
|
||||
t = ctx.timeout("cmake-cpu")
|
||||
t0 = time.time()
|
||||
for label, cmd in (("configure", [cmake, "--preset", "cpu-check", "-B", build]),
|
||||
("build", [cmake, "--build", build, "--parallel", str(ctx.jobs)])):
|
||||
rc, _ = run_cmd(cmd, ROOT, env, max(30, t - (time.time() - t0)), log)
|
||||
if rc is None or rc != 0:
|
||||
r.seconds = time.time() - t0
|
||||
r.failures.append("cmake " + label + (" (timed out)" if rc is None else ""))
|
||||
r.reason = " | ".join(l.strip() for l in tail(log, 4))[:300]
|
||||
return finish(r, rc if rc is not None else 1, rc is None, log)
|
||||
junit = ctx.scratch.path("ctest.xml")
|
||||
rc, _ = run_cmd([shutil.which("ctest"), "--test-dir", build, "--output-on-failure", "--parallel", str(min(ctx.jobs, 8)),
|
||||
"--timeout", "120", "--no-tests=error", "--output-junit", junit], ROOT, env, max(30, t - (time.time() - t0)), log)
|
||||
r.seconds = time.time() - t0
|
||||
if rc is None:
|
||||
return finish(r, 1, True, log)
|
||||
if not parse_junit(junit, r):
|
||||
r.failures.append("ctest wrote no result file")
|
||||
return finish(r, 1, False, log)
|
||||
# the one-line summary the CMakeLists prints at configure time (which kinds of test were registered)
|
||||
for line in open(log, "rb").read().decode("utf-8", "replace").splitlines():
|
||||
if "strata tests:" in line:
|
||||
r.reason = line.split("-- ", 1)[-1].strip()
|
||||
break
|
||||
return finish(r, rc, False, log)
|
||||
|
||||
|
||||
def step_pytest(name, targets, extra=(), need_node=False, cwd=None):
|
||||
def run(ctx):
|
||||
r = Result(name)
|
||||
if not have_module("pytest"):
|
||||
r.reason = "pytest is not installed for %s" % sys.executable
|
||||
return r
|
||||
if need_node:
|
||||
node, why = node_version()
|
||||
if not node:
|
||||
r.reason = why
|
||||
return r
|
||||
paths = []
|
||||
for t in targets:
|
||||
hit = sorted(glob.glob(os.path.join(ROOT, t))) if any(c in t for c in "*?[") else ([os.path.join(ROOT, t)] if os.path.exists(os.path.join(ROOT, t)) else [])
|
||||
paths.extend(hit)
|
||||
if not paths:
|
||||
r.reason = "no test files found for " + " ".join(targets)
|
||||
return r
|
||||
if ctx.shard: # --shard I/N: every N-th test file (for a machine that allows only short jobs); the shards together are the whole step
|
||||
i, n = ctx.shard
|
||||
files = []
|
||||
for pth in paths:
|
||||
if os.path.isdir(pth):
|
||||
for pat in ("test_*.py", "*_test.py"):
|
||||
files += glob.glob(os.path.join(pth, "**", pat), recursive=True)
|
||||
else:
|
||||
files.append(pth)
|
||||
skip_dirs = [os.path.normpath(x[len("--ignore="):]) for x in extra if x.startswith("--ignore=")]
|
||||
files = sorted({f for f in files if not any(os.path.normpath(f).startswith(d + os.sep) for d in skip_dirs)})
|
||||
paths = files[i - 1::n]
|
||||
r.reason = "shard %d/%d: %d of %d test files" % (i, n, len(paths), len(files))
|
||||
if not paths:
|
||||
return r
|
||||
junit = ctx.scratch.path(name + ".xml")
|
||||
log = ctx.scratch.path(name + ".log")
|
||||
cmd = [sys.executable, "-m", "pytest", "-q", "-p", "no:cacheprovider", "-o", "junit_family=xunit2",
|
||||
"--junitxml=" + junit, "--maxfail=1000000", *extra, *paths]
|
||||
rc, secs = run_cmd(cmd, cwd or ROOT, ctx.scratch.env(), ctx.timeout(name), log)
|
||||
r.seconds = secs
|
||||
ok = parse_junit(junit, r)
|
||||
if rc is None and not ok:
|
||||
return finish(r, 1, True, log)
|
||||
if not ok and rc not in (0, 5):
|
||||
r.failures.append("pytest wrote no result file (exit %s)" % rc)
|
||||
r.reason = " | ".join(l.strip() for l in tail(log, 4))[:300]
|
||||
if rc == 5: # nothing collected
|
||||
r.reason = "pytest collected no tests"
|
||||
r.status = "not run"
|
||||
return r
|
||||
return finish(r, rc, rc is None, log)
|
||||
return run
|
||||
|
||||
|
||||
def step_optional(ctx):
|
||||
r = Result("py-optional")
|
||||
for mod, what in (("torch", "torch"), ("gguf", "gguf (llama.cpp's gguf-py: pip install gguf)")):
|
||||
if not have_module(mod):
|
||||
r.reason = "%s is not installed for %s (tools/dflash2 and tools/dsv41 tests need torch and gguf)" % (what, sys.executable)
|
||||
return r
|
||||
return step_pytest("py-optional", ["tools/dflash2", "tools/dsv41", "tools/trellis"])(ctx)
|
||||
|
||||
|
||||
def npm_cmd():
|
||||
return shutil.which("npm.cmd" if WIN else "npm") or shutil.which("npm")
|
||||
|
||||
|
||||
def host_ready(ctx, r):
|
||||
"""Node, npm and apps/host/node_modules. Installs the packages with `npm ci` when they are missing."""
|
||||
node, why = node_version()
|
||||
if not node:
|
||||
r.reason = why
|
||||
return False
|
||||
if not npm_cmd():
|
||||
r.reason = "npm is not installed"
|
||||
return False
|
||||
if os.path.isdir(os.path.join(HOST, "node_modules", ".bin")) or ctx.npm_ready:
|
||||
return True
|
||||
log = ctx.scratch.path("npm-ci.log")
|
||||
rc, _ = run_cmd([npm_cmd(), "ci", "--no-audit", "--no-fund"], HOST, ctx.scratch.env(), 900, log)
|
||||
if rc != 0:
|
||||
r.reason = "npm ci could not install apps/host's packages (no network?); log " + log
|
||||
return False
|
||||
ctx.npm_ready = True
|
||||
return True
|
||||
|
||||
|
||||
def step_npm(name, script, args=()):
|
||||
def run(ctx):
|
||||
r = Result(name)
|
||||
if not host_ready(ctx, r):
|
||||
return r
|
||||
log = ctx.scratch.path(name + ".log")
|
||||
junit = ctx.scratch.path(name + ".xml")
|
||||
cmd = [npm_cmd(), "run", script]
|
||||
if name == "host-test":
|
||||
# the junit file gives the counts and the skip reasons; spec goes to the log
|
||||
cmd += ["--", "--test-reporter=spec", "--test-reporter-destination=stdout", "--test-reporter=junit",
|
||||
"--test-reporter-destination=" + junit]
|
||||
rc, secs = run_cmd(cmd, HOST, ctx.scratch.env({"STRATA_TEST_TIMEOUT_MS": "180000"}), ctx.timeout(name), log)
|
||||
r.seconds = secs
|
||||
if name == "host-test" and os.path.exists(junit):
|
||||
parse_junit(junit, r)
|
||||
# node writes suites and tests as testcases; a suite-level failure has no testcase of its own
|
||||
elif name == "host-test":
|
||||
summary = {m.group(1): int(m.group(2)) for m in re.finditer(r"^\W*(tests|pass|fail|cancelled|skipped)\s+(\d+)\s*$",
|
||||
open(log, "rb").read().decode("utf-8", "replace"), re.M)}
|
||||
r.passed, r.failed, r.skipped = summary.get("pass", 0), summary.get("fail", 0) + summary.get("cancelled", 0), summary.get("skipped", 0)
|
||||
if r.skipped:
|
||||
r.skips["(node:test gives no reason in this report)"] = r.skipped
|
||||
return finish(r, rc, rc is None, log)
|
||||
return run
|
||||
|
||||
|
||||
STEP_FUNCS = {
|
||||
"cmake-cpu": step_cmake,
|
||||
"py-tests": step_pytest("py-tests", ["tests"], ["--ignore=" + os.path.join(ROOT, "tests", "snapshot")]),
|
||||
"py-serve": step_pytest("py-serve", ["serve"]),
|
||||
"py-tools": step_pytest("py-tools", ["tools/test_*.py"]),
|
||||
"py-optional": step_optional,
|
||||
"host-typecheck": step_npm("host-typecheck", "typecheck"),
|
||||
"host-test": step_npm("host-test", "test"),
|
||||
"host-build": step_npm("host-build", "build"),
|
||||
"snapshot": step_pytest("snapshot", ["tests/snapshot"], need_node=True),
|
||||
}
|
||||
|
||||
|
||||
class Ctx:
|
||||
def __init__(self, scratch, scale, jobs, shard=None):
|
||||
self.scratch, self.scale, self.jobs, self.npm_ready, self.shard = scratch, scale, jobs, False, shard
|
||||
|
||||
def timeout(self, name):
|
||||
return int(TIMEOUTS[name] * self.scale)
|
||||
|
||||
|
||||
def table(results):
|
||||
head = ("step", "status", "passed", "failed", "skipped", "time", "")
|
||||
rows = [head] + [(r.name, r.status.upper() if r.status != "passed" else "passed", str(r.passed), str(r.failed), str(r.skipped),
|
||||
"%dm%02ds" % divmod(int(r.seconds), 60), r.reason if r.status != "passed" else "") for r in results]
|
||||
w = [max(len(row[i]) for row in rows) for i in range(6)]
|
||||
out = []
|
||||
for i, row in enumerate(rows):
|
||||
out.append(" ".join(row[j].ljust(w[j]) if j in (0, 1) else row[j].rjust(w[j]) for j in range(6)) + (" " + row[6] if row[6] else ""))
|
||||
if i == 0:
|
||||
out.append(" ".join("-" * x for x in w))
|
||||
return "\n".join(out)
|
||||
|
||||
|
||||
def main(argv=None):
|
||||
ap = argparse.ArgumentParser(description="Strata fast checks (see the top of this file)")
|
||||
ap.add_argument("--only", help="comma list of steps")
|
||||
ap.add_argument("--skip", help="comma list of steps to leave out")
|
||||
ap.add_argument("--list", action="store_true", help="list the steps and their timeouts")
|
||||
ap.add_argument("--scratch", help="scratch folder (default: a new one under the temp folder, removed at the end)")
|
||||
ap.add_argument("--keep", action="store_true", help="keep the scratch folder (logs, build tree)")
|
||||
ap.add_argument("--json", help="also write the result to this file")
|
||||
ap.add_argument("--strict", action="store_true", help="exit 2 when a step could not run")
|
||||
ap.add_argument("--timeout-scale", type=float, default=1.0, help="multiply every step's timeout")
|
||||
ap.add_argument("--shard", help="I/N: run every N-th test file of the Python steps (e.g. 1/4), for a machine with short job limits")
|
||||
ap.add_argument("--jobs", type=int, default=max(2, min(16, (os.cpu_count() or 4) // 2)), help="build/test parallelism")
|
||||
a = ap.parse_args(argv)
|
||||
if a.list:
|
||||
for s in STEPS:
|
||||
print("%-15s timeout %ds" % (s, TIMEOUTS[s]))
|
||||
return 0
|
||||
want = [s.strip() for s in a.only.split(",")] if a.only else list(STEPS)
|
||||
for s in want + [x.strip() for x in (a.skip or "").split(",") if x.strip()]:
|
||||
if s not in STEPS:
|
||||
ap.error("unknown step %r (try --list)" % s)
|
||||
want = [s for s in want if s not in [x.strip() for x in (a.skip or "").split(",")]]
|
||||
scratch = Scratch(a.scratch, a.keep)
|
||||
shard = None
|
||||
if a.shard:
|
||||
m = re.match(r"^(\d+)/(\d+)$", a.shard)
|
||||
if not m or not 1 <= int(m.group(1)) <= int(m.group(2)):
|
||||
ap.error("--shard wants I/N with 1 <= I <= N, for example 2/4")
|
||||
shard = (int(m.group(1)), int(m.group(2)))
|
||||
ctx = Ctx(scratch, a.timeout_scale, a.jobs, shard)
|
||||
results = []
|
||||
print("tools/check: %s, python %s, scratch %s" % (time.strftime("%Y-%m-%d %H:%M:%S"), sys.version.split()[0], scratch.dir), flush=True)
|
||||
try:
|
||||
for s in want:
|
||||
print("... %s" % s, flush=True)
|
||||
try:
|
||||
r = STEP_FUNCS[s](ctx)
|
||||
except Exception as e: # noqa: BLE001 a bug in this script is a failure, never a pass
|
||||
r = Result(s)
|
||||
r.status, r.reason = "failed", "tools/check itself failed: %r" % (e,)
|
||||
r.failures.append(r.reason)
|
||||
print(" %s %s" % (s, r.status), flush=True)
|
||||
results.append(r)
|
||||
print()
|
||||
print(table(results))
|
||||
for r in results:
|
||||
if r.failures:
|
||||
print("\nFAILED in %s (%d):" % (r.name, len(r.failures)))
|
||||
for f in r.failures[:40]:
|
||||
print(" " + f)
|
||||
if len(r.failures) > 40:
|
||||
print(" ... %d more (see %s)" % (len(r.failures) - 40, r.log))
|
||||
if r.log:
|
||||
print(" log: " + r.log)
|
||||
skips = [(r.name, k, v) for r in results for k, v in sorted(r.skips.items(), key=lambda kv: -kv[1])]
|
||||
if skips:
|
||||
print("\nSKIPPED, with the reason:")
|
||||
for name, reason, n in skips:
|
||||
print(" %-14s %5d %s" % (name, n, reason))
|
||||
nrun = [r for r in results if r.status == "not run"]
|
||||
failed = [r for r in results if r.status == "failed"]
|
||||
print("\ncheck: %d passed, %d failed, %d not run (of %d steps)" % (
|
||||
len([r for r in results if r.status == "passed"]), len(failed), len(nrun), len(results)))
|
||||
if a.json:
|
||||
with open(a.json, "w", encoding="utf-8") as f:
|
||||
json.dump([r.row() for r in results], f, indent=1)
|
||||
if a.keep or a.scratch:
|
||||
print("scratch kept: " + scratch.dir)
|
||||
return 1 if failed else (2 if (a.strict and nrun) else 0)
|
||||
finally:
|
||||
scratch.cleanup()
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
sys.exit(main())
|
||||
@@ -3,7 +3,7 @@
|
||||
Everything is fp32 on the CPU. The target's embedding table and LM head are the dequantized tensors of OUR GGUF (IQ4_XS embedding,
|
||||
Q6_K head), the context features are the hidden states dumped by the engine's STRATA_TAP_FILE hook.
|
||||
"""
|
||||
import json, math, struct, sys
|
||||
import json, struct, sys
|
||||
import numpy as np
|
||||
import torch
|
||||
import torch.nn.functional as F
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
"""Does a Q8_0 drafter (the matrices block-quantized like ggml Q8_0) keep the acceptance? Walk over two prompts, block 4 and 8."""
|
||||
import json, sys, time
|
||||
import json, sys
|
||||
from pathlib import Path
|
||||
import numpy as np
|
||||
import torch
|
||||
|
||||
@@ -1,3 +1,13 @@
|
||||
# A lab script, not a unit test: it needs the owner PC's work folder (the drafter's weights and config) and the reference model
|
||||
# file under phaseA-tests/ornith_tmp. pytest collects it as a module, so it says so and skips when they are not there.
|
||||
import os
|
||||
_LAB = (r"C:\dflash2-work", r"C:\Users\AI-Server\Desktop\Strata\phaseA-tests\ornith_tmp\gh_model.py")
|
||||
if not all(os.path.exists(p) for p in _LAB):
|
||||
try:
|
||||
import pytest
|
||||
pytest.skip("needs the DFlash2 lab files on the owner PC: " + " and ".join(_LAB), allow_module_level=True)
|
||||
except ImportError:
|
||||
raise SystemExit("needs the DFlash2 lab files on the owner PC: " + " and ".join(_LAB))
|
||||
import sys, importlib.util, json, torch
|
||||
sys.path.insert(0, r"C:\dflash2-work")
|
||||
import dflash2_ref as R
|
||||
|
||||
@@ -5,7 +5,6 @@ two configs and compare the answers.
|
||||
python quality_fused.py strata-iq3s-nofused.json strata-iq3s-best.json out.json
|
||||
"""
|
||||
import json
|
||||
import os
|
||||
import subprocess
|
||||
import sys
|
||||
import time
|
||||
|
||||
@@ -17,7 +17,6 @@ import subprocess
|
||||
import sys
|
||||
import tempfile
|
||||
import time
|
||||
import urllib.request
|
||||
from pathlib import Path
|
||||
|
||||
ROOT = Path(__file__).resolve().parents[1]
|
||||
|
||||
@@ -15,9 +15,7 @@ import ctypes
|
||||
import json
|
||||
import multiprocessing as mp
|
||||
import os
|
||||
import pathlib
|
||||
import re
|
||||
import sys
|
||||
import time
|
||||
|
||||
import numpy as np
|
||||
|
||||
@@ -19,7 +19,6 @@ from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import json
|
||||
import math
|
||||
import pathlib
|
||||
import struct
|
||||
import sys
|
||||
|
||||
@@ -23,11 +23,8 @@ from __future__ import annotations
|
||||
|
||||
import argparse
|
||||
import hashlib
|
||||
import json
|
||||
import os
|
||||
import pathlib
|
||||
import re
|
||||
import shutil
|
||||
import subprocess
|
||||
import sys
|
||||
import time
|
||||
|
||||
@@ -19,7 +19,6 @@ import json
|
||||
import math
|
||||
import os
|
||||
import struct
|
||||
import sys
|
||||
|
||||
|
||||
WARM = 0 # accesses before this index warm the cache but are not counted
|
||||
|
||||
@@ -8,10 +8,7 @@ hottest experts in (the engine's preload) and the same cold-start window's hit r
|
||||
(profile only, no adaptation)."""
|
||||
import argparse
|
||||
import collections
|
||||
import io
|
||||
import json
|
||||
import struct
|
||||
import sys
|
||||
|
||||
import numpy as np
|
||||
|
||||
|
||||
@@ -29,7 +29,7 @@ Rules (paper): mismatch = the scorer's argmax is not the forced token; margin =
|
||||
mismatch with margin > 2 delta. REJECT if the candidate has more events than stock by an exact one-sided binomial test (p < 0.05,
|
||||
exposure-weighted). (The paper also audits events stock reproduces on its own paths; that audit is not done here, events are counted raw.)
|
||||
"""
|
||||
import argparse, json, math, os, random, signal, subprocess, sys, time, pathlib
|
||||
import argparse, json, math, os, random, signal, subprocess, sys, pathlib
|
||||
|
||||
SCORER_DROP_VALUE = {"--draft", "--draft-model", "--draft-fallback", "--mtp-draft-vocab", "--spec", "--draft-conf", "--draft-cap", "--mtp", "--spec-min-p",
|
||||
"--prompt-cache", "--max-context", "--prefill", "--short-read", "--prompt-cache-every", "--conversation-cache-mib", "--suffix-draft"}
|
||||
|
||||
@@ -25,7 +25,6 @@ from __future__ import annotations
|
||||
|
||||
import json
|
||||
import pathlib
|
||||
import struct
|
||||
import sys
|
||||
|
||||
import numpy as np
|
||||
|
||||
+1
-1
@@ -1,6 +1,6 @@
|
||||
#!/usr/bin/env python3
|
||||
"""recon.py - reconcile PAR-A's boston numbers with this harness (runs on boston, one arm after the other)."""
|
||||
import json, os, subprocess, sys, time
|
||||
import os, subprocess, sys, time
|
||||
PA = "/data/external/niko/parA"
|
||||
I = "/data/external/niko/int029"
|
||||
R = "/data/external/niko/par29"
|
||||
|
||||
Loaded 100 of 126 files, more files were not shown because too many files have changed in this diff.
Show more
Reference in new issue
Block a user