wehub-resource-sync
426e9eeabd
Voice Workbench / headless workbench (mocked backends) (push) Has been cancelled
Voice Workbench / real acoustic lane (nightly, provisioned only) (push) Has been cancelled
ci / test (push) Has been cancelled
ci / lint-and-format (push) Has been cancelled
ci / build (push) Has been cancelled
ci / dev-startup (push) Has been cancelled
gitleaks / gitleaks (push) Has been cancelled
Markdown Links / Relative Markdown Links (push) Has been cancelled
Quality (Extended) / Homepage Build (PR smoke) (push) Has been cancelled
Quality (Extended) / Comment-only diff guard (push) Has been cancelled
Quality (Extended) / Format + Type Safety Ratchet (push) Has been cancelled
Quality (Extended) / Develop Gate (secret scan + UI determinism) (push) Has been cancelled
Quality (Extended) / Develop Gate (lint) (push) Has been cancelled
Chat shell gestures / Chat shell gesture + parity e2e (push) Has been cancelled
Cloud Gateway Discord / Test (push) Has been cancelled
Benchmark Bridge Tests / benchmark (bunx @biomejs/biome check packages/lifeops-bench/src, benchmark-lint) (push) Has been cancelled
Benchmark Bridge Tests / benchmark (bunx vitest run --config packages/lifeops-bench/vitest.config.ts --root packages/lifeops-bench --passWithNoTests, benchmark-tests) (push) Has been cancelled
Build Agent Image / build-and-push (push) Has been cancelled
Dev Smoke / bun run dev onboarding chat (push) Has been cancelled
Dev Smoke / Vite HMR dependency-level smoke (push) Has been cancelled
Electrobun Submodule Guard / electrobun gitlink is fetchable (push) Has been cancelled
Publish @elizaos/example-code / check_npm (push) Has been cancelled
Publish @elizaos/example-code / publish_npm (push) Has been cancelled
Publish @elizaos/plugin-elizacloud / verify_version (push) Has been cancelled
Publish @elizaos/plugin-elizacloud / publish_npm (push) Has been cancelled
Sandbox Live Smoke / Sandbox live smoke (push) Has been cancelled
Snap Build & Test / Build Snap (amd64) (push) Has been cancelled
Snap Build & Test / Build Snap (arm64) (push) Has been cancelled
Test Packaging / elizaos CLI global-install smoke (node + bun) (push) Has been cancelled
Cloud Gateway Webhook / Test (push) Has been cancelled
Cloud Tests / lint-and-types (push) Has been cancelled
Cloud Tests / unit-tests (push) Has been cancelled
Cloud Tests / integration-tests (push) Has been cancelled
Cloud Tests / e2e-tests (push) Has been cancelled
CodeQL Advanced / Analyze (javascript-typescript) (push) Has been cancelled
Deploy Apps Worker (Product 2) / Determine environment (push) Has been cancelled
Deploy Apps Worker (Product 2) / Deploy apps worker to apps-control host (${{ needs.determine-env.outputs.environment }}) (push) Has been cancelled
Deploy Eliza Provisioning Worker / Determine environment (push) Has been cancelled
Deploy Eliza Provisioning Worker / Deploy worker to Hetzner host (${{ needs.determine-env.outputs.environment }} @ ${{ needs.determine-env.outputs.deployment_sha }}) (push) Has been cancelled
Dev Smoke / Classify changed paths (push) Has been cancelled
supply-chain / sbom (push) Has been cancelled
supply-chain / vulnerability-scan (push) Has been cancelled
Build, Push & Deploy to Phala Cloud / build-and-push (push) Has been cancelled
Test Packaging / Validate Packaging Configs (push) Has been cancelled
Test Packaging / Build & Test PyPI Package (push) Has been cancelled
Test Packaging / PyPI on Python ${{ matrix.python }} (push) Has been cancelled
Test Packaging / Pack & Test JS Tarballs (push) Has been cancelled
UI Fixture E2E / ui-fixture-e2e (push) Has been cancelled
UI Fixture E2E / fixture-e2e (push) Has been cancelled
UI Story Gate / story-gate (push) Has been cancelled
vault-ci / test (macos-latest) (push) Has been cancelled
vault-ci / test (ubuntu-latest) (push) Has been cancelled
vault-ci / test (windows-latest) (push) Has been cancelled
vault-ci / app-core wiring tests (push) Has been cancelled
verify-patches / verify patches/CHECKSUMS.sha256 (push) Has been cancelled
Voice Benchmark Smoke / voice-emotion fixture smoke (push) Has been cancelled
Voice Benchmark Smoke / voiceagentbench fixture smoke (push) Has been cancelled
Voice Benchmark Smoke / voicebench-quality unit smoke (push) Has been cancelled
Voice Benchmark Smoke / voicebench TypeScript unit (no audio) (push) Has been cancelled
Voice Benchmark Smoke / voice bench smoke summary (push) Has been cancelled
Windows CI / windows ([bun run --cwd packages/app-core test bun run --cwd packages/elizaos test bun run --cwd packages/cloud/shared test], app-and-cli) (push) Has been cancelled
Windows CI / windows ([bun run --cwd packages/scenario-runner test bun run --cwd packages/vault test bun run --cwd packages/security test bun run --cwd plugins/plugin-coding-tools test], framework-packages) (push) Has been cancelled
Windows CI / windows ([bun run --cwd plugins/plugin-elizacloud test bun run --cwd plugins/plugin-discord test bun run --cwd plugins/plugin-anthropic test bun run --cwd plugins/plugin-openai test bun run --cwd plugins/plugin-app-control test bun run --cwd plugins/pl… (push) Has been cancelled
Windows CI / windows ([node packages/scripts/run-turbo.mjs run build --filter=@elizaos/core --filter=@elizaos/shared --filter=@elizaos/agent --concurrency=4 node packages/scripts/run-bash-linux-only.mjs scripts/verify-riscv64-buildpaths.sh node packages/scripts/run… (push) Has been cancelled
Windows CI / windows ([node packages/scripts/run-turbo.mjs run typecheck --filter=@elizaos/core --filter=@elizaos/shared --filter=@elizaos/cloud-shared --concurrency=4 bun run --cwd packages/core test bun run --cwd packages/shared test], core-runtime, 75) (push) Has been cancelled
955 lines
57 KiB
Makefile
955 lines
57 KiB
Makefile
# DRAFT verification Makefile for the turbo3 / turbo4 / turbo3_tcq kernels.
|
||
# Builds the platform-specific harnesses only when the relevant SDK is present.
|
||
# This file MUST stay no-op on plain Linux without a GPU SDK so that CI is
|
||
# safe to run.
|
||
|
||
CC ?= cc
|
||
CXX ?= c++
|
||
CFLAGS ?= -O2 -Wall -Wextra -std=c11
|
||
CXXFLAGS?= -O2 -Wall -Wextra -std=c++17
|
||
|
||
REF_SRC = ../reference/turbo_kernels.c
|
||
REF_OBJ = turbo_kernels.o
|
||
QJL_POLAR_SRC = qjl_polar_ref.c
|
||
QJL_POLAR_OBJ = qjl_polar_ref.o
|
||
|
||
.PHONY: all reference-test kernel-contract istft-verify two-agent-voice-demo vulkan vulkan-spirv vulkan-verify vulkan-verify-multiblock vulkan-verify-fallbacks vulkan-verify-fused vulkan-qjl-polar vulkan-bench vulkan-dispatch-smoke vulkan-native-smoke android-vulkan-smoke android-vulkan-graph-smoke metal metal-verify metal-verify-shipped metal-verify-multiblock metal-verify-fused metal-bench metal-bench-batched metal-bench-multiblock metal-bench-fused metal-upscale-probe metal-upscale-probe-require cpu-dispatch-smoke cpu-qjl-polar-attn-smoke cpu-bench bench cuda cuda-verify cuda-verify-fused cuda-hardware rocm-hardware gh200-hardware windows-hardware verify-fork dispatch-smoke dispatch-smoke-implemented mmproj-verify force clean
|
||
|
||
all: reference-test
|
||
@echo "[kernels/verify] Built reference test."
|
||
@echo "[kernels/verify] To build the Vulkan harness: VULKAN_SDK=/path/to/sdk make vulkan"
|
||
@echo "[kernels/verify] To build the Metal harness (macOS): make metal"
|
||
|
||
$(REF_OBJ): $(REF_SRC) ../reference/turbo_kernels.h
|
||
$(CC) $(CFLAGS) -c $< -o $@
|
||
|
||
$(QJL_POLAR_OBJ): $(QJL_POLAR_SRC) qjl_polar_ref.h ../reference/turbo_kernels.h
|
||
$(CC) $(CFLAGS) -c $< -o $@
|
||
|
||
reference-test: gen_fixture
|
||
@echo "[kernels/verify] reference vs reference (fixture sanity)..."
|
||
./gen_fixture --self-test
|
||
|
||
kernel-contract:
|
||
node check_kernel_contract.mjs
|
||
|
||
istft-verify:
|
||
node istft-verify.mjs --backend cpu --tol 1e-3
|
||
|
||
two-agent-voice-demo:
|
||
node two_agent_voice_demo.mjs --backend synthetic --turns 4 --json
|
||
|
||
gen_fixture: gen_fixture.c $(REF_OBJ) $(QJL_POLAR_OBJ)
|
||
$(CC) $(CFLAGS) -I../reference $< $(REF_OBJ) $(QJL_POLAR_OBJ) -lm -o $@
|
||
|
||
# Vulkan harness. Resolution order:
|
||
# 1. VULKAN_SDK env var (LunarG SDK install) — preferred, gives glslangValidator
|
||
# 2. Fetched Khronos Vulkan-Headers checkout from the mtp build script
|
||
# (~/.cache/eliza-mtp/vulkan-headers/include) + system libvulkan.so.1
|
||
# 3. Fail with a clear instruction.
|
||
#
|
||
# Path 2 lets the harness build on stock Linux machines that already have
|
||
# `bun run packages/app-core/scripts/build-llama-cpp-mtp.mjs --target
|
||
# linux-x64-vulkan` set up but no Vulkan SDK installed.
|
||
ELIZA_MTP_VULKAN_HEADERS_DIR ?= $(HOME)/.cache/eliza-mtp/vulkan-headers/include
|
||
SYSTEM_LIBVULKAN := $(firstword \
|
||
$(wildcard /usr/lib/x86_64-linux-gnu/libvulkan.so.1) \
|
||
$(wildcard /usr/lib64/libvulkan.so.1) \
|
||
$(wildcard /usr/lib/libvulkan.so.1))
|
||
|
||
# Homebrew Vulkan SDK on macOS (vulkan-headers + vulkan-loader + molten-vk).
|
||
# Detected by the loader dylib + headers under /opt/homebrew.
|
||
HOMEBREW_VULKAN_HEADERS := $(wildcard /opt/homebrew/include/vulkan/vulkan.h)
|
||
HOMEBREW_VULKAN_LIB := $(wildcard /opt/homebrew/opt/vulkan-loader/lib/libvulkan.dylib)
|
||
|
||
vulkan: vulkan_verify
|
||
vulkan_verify: vulkan_verify.cpp $(REF_OBJ) $(QJL_POLAR_OBJ)
|
||
ifdef VULKAN_SDK
|
||
$(CXX) $(CXXFLAGS) -I$(VULKAN_SDK)/include -I../reference -I. \
|
||
vulkan_verify.cpp $(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
-L$(VULKAN_SDK)/lib -lvulkan -lm -o vulkan_verify
|
||
else ifneq ($(strip $(HOMEBREW_VULKAN_HEADERS)$(HOMEBREW_VULKAN_LIB)),)
|
||
$(CXX) $(CXXFLAGS) -I/opt/homebrew/include -I../reference -I. \
|
||
vulkan_verify.cpp $(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
-L/opt/homebrew/opt/vulkan-loader/lib -lvulkan -lm -o vulkan_verify
|
||
@echo "[vulkan_verify] Built against Homebrew vulkan-loader (macOS / MoltenVK)."
|
||
@echo "[vulkan_verify] Run with:"
|
||
@echo " export VK_ICD_FILENAMES=/opt/homebrew/etc/vulkan/icd.d/MoltenVK_icd.json"
|
||
@echo " export DYLD_LIBRARY_PATH=/opt/homebrew/opt/vulkan-loader/lib:/opt/homebrew/opt/molten-vk/lib"
|
||
else
|
||
ifeq ($(wildcard $(ELIZA_MTP_VULKAN_HEADERS_DIR)/vulkan/vulkan.hpp),)
|
||
@echo "Vulkan harness needs one of:"
|
||
@echo " - VULKAN_SDK pointing at LunarG, or"
|
||
@echo " - macOS: brew install molten-vk vulkan-headers vulkan-loader, or"
|
||
@echo " - Fetched Khronos headers at $(ELIZA_MTP_VULKAN_HEADERS_DIR)"
|
||
@echo " (run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target linux-x64-vulkan' to fetch them)"
|
||
@false
|
||
else ifeq ($(SYSTEM_LIBVULKAN),)
|
||
@echo "Vulkan harness needs libvulkan.so.1 (apt install libvulkan1)"
|
||
@false
|
||
else
|
||
$(CXX) $(CXXFLAGS) -I$(ELIZA_MTP_VULKAN_HEADERS_DIR) -I../reference -I. \
|
||
vulkan_verify.cpp $(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
$(SYSTEM_LIBVULKAN) -lm -o vulkan_verify
|
||
endif
|
||
endif
|
||
|
||
# vulkan_bench — GPU-time benchmark harness (VK_QUERY_TYPE_TIMESTAMP). Sibling
|
||
# of vulkan_verify; measures per-dispatch GPU time and sweeps the multi-block
|
||
# spec constant. Self-contained (only the Vulkan headers + libvulkan), so the
|
||
# header/lib resolution mirrors vulkan_verify but it does not link the C
|
||
# reference objects.
|
||
vulkan_bench: vulkan_bench.cpp
|
||
ifdef VULKAN_SDK
|
||
$(CXX) $(CXXFLAGS) -I$(VULKAN_SDK)/include -I. vulkan_bench.cpp \
|
||
-L$(VULKAN_SDK)/lib -lvulkan -lm -o vulkan_bench
|
||
else ifneq ($(strip $(HOMEBREW_VULKAN_HEADERS)$(HOMEBREW_VULKAN_LIB)),)
|
||
$(CXX) $(CXXFLAGS) -I/opt/homebrew/include -I. vulkan_bench.cpp \
|
||
-L/opt/homebrew/opt/vulkan-loader/lib -lvulkan -lm -o vulkan_bench
|
||
else
|
||
ifeq ($(wildcard $(ELIZA_MTP_VULKAN_HEADERS_DIR)/vulkan/vulkan.hpp),)
|
||
@echo "vulkan_bench needs Vulkan headers (VULKAN_SDK, brew vulkan-headers, or the fetched Khronos headers)."
|
||
@false
|
||
else ifeq ($(SYSTEM_LIBVULKAN),)
|
||
@echo "vulkan_bench needs libvulkan.so.1 (apt install libvulkan1)"
|
||
@false
|
||
else
|
||
$(CXX) $(CXXFLAGS) -I$(ELIZA_MTP_VULKAN_HEADERS_DIR) -I. vulkan_bench.cpp \
|
||
$(SYSTEM_LIBVULKAN) -lm -o vulkan_bench
|
||
endif
|
||
endif
|
||
|
||
# Run the Vulkan kernel benchmark on the local device. Disables the Mesa shader
|
||
# cache for clean timings; refuses software ICDs (lavapipe) unless
|
||
# ELIZA_ALLOW_SOFTWARE_VULKAN=1. Writes JSON to VULKAN_BENCH_JSON if set.
|
||
VULKAN_BENCH_JSON ?=
|
||
vulkan-bench: vulkan_bench vulkan-spirv
|
||
MESA_SHADER_CACHE_DISABLE=1 ./vulkan_bench $(if $(VULKAN_BENCH_JSON),--json $(VULKAN_BENCH_JSON),)
|
||
|
||
# mul_mm GEMM micro-benchmark (#9584/#9715): scalar vs shared-mem-tiled vs
|
||
# register-blocked, the cooperative-matrix-free prefill mul_mat study. Same
|
||
# header/lib resolution as vulkan_bench. On Mali-G715 the register-blocked
|
||
# variant is the only one that beats scalar (~1.4-2.3x); naive tiling loses.
|
||
mul_mm_bench: mul_mm_bench.cpp
|
||
ifdef VULKAN_SDK
|
||
$(CXX) $(CXXFLAGS) -I$(VULKAN_SDK)/include -I. mul_mm_bench.cpp -L$(VULKAN_SDK)/lib -lvulkan -lm -o mul_mm_bench
|
||
else ifneq ($(strip $(HOMEBREW_VULKAN_HEADERS)$(HOMEBREW_VULKAN_LIB)),)
|
||
$(CXX) $(CXXFLAGS) -I/opt/homebrew/include -I. mul_mm_bench.cpp -L/opt/homebrew/opt/vulkan-loader/lib -lvulkan -lm -o mul_mm_bench
|
||
else ifneq ($(SYSTEM_LIBVULKAN),)
|
||
$(CXX) $(CXXFLAGS) -I$(ELIZA_MTP_VULKAN_HEADERS_DIR) -I. mul_mm_bench.cpp $(SYSTEM_LIBVULKAN) -lm -o mul_mm_bench
|
||
else
|
||
@echo "mul_mm_bench needs Vulkan headers + libvulkan (VULKAN_SDK / brew / system)."; false
|
||
endif
|
||
|
||
mul_mm-bench: mul_mm_bench vulkan-spirv
|
||
MESA_SHADER_CACHE_DISABLE=1 ./mul_mm_bench --spv-dir ../vulkan
|
||
|
||
# Run vulkan_verify against the QJL + Polar fixtures (needs a Vulkan-capable
|
||
# host — fails cleanly with "no Vulkan devices" on machines without an ICD).
|
||
# Requires the SPIR-V artifacts under ../vulkan/*.spv: build them with
|
||
# `make vulkan-spirv` first. Pre-supposes that gen_fixture has produced
|
||
# fixtures/qjl.json + fixtures/polar.json (run `make reference-test` first).
|
||
vulkan-qjl-polar: vulkan_verify vulkan-spirv
|
||
@echo "[kernels/verify] running vulkan_verify on QJL fixture..."
|
||
./vulkan_verify ../vulkan/qjl.spv fixtures/qjl.json
|
||
@echo "[kernels/verify] running vulkan_verify on Polar fixture..."
|
||
./vulkan_verify ../vulkan/polar.spv fixtures/polar.json
|
||
@echo "[kernels/verify] running vulkan_verify on Polar+QJL-residual fixture..."
|
||
./vulkan_verify ../vulkan/polar.spv fixtures/polar_qjl.json
|
||
@echo "[kernels/verify] running vulkan_verify on Polar pre-Hadamard-query fixture..."
|
||
./vulkan_verify ../vulkan/polar_preht.spv fixtures/polar.json
|
||
@echo "[kernels/verify] running vulkan_verify on Polar pre-Hadamard-query + QJL-residual fixture..."
|
||
./vulkan_verify ../vulkan/polar_preht.spv fixtures/polar_qjl.json
|
||
|
||
# Full Vulkan correctness gate for the default-on kernel set. Runs all five
|
||
# score kernels plus Polar residual and pre-Hadamard variants so the optional
|
||
# residual branch and H*q hot path are not left to source review only.
|
||
vulkan-verify: vulkan_verify vulkan-spirv reference-test
|
||
@echo "[kernels/verify] vulkan turbo3..."
|
||
./vulkan_verify ../vulkan/turbo3.spv fixtures/turbo3.json
|
||
@echo "[kernels/verify] vulkan turbo4..."
|
||
./vulkan_verify ../vulkan/turbo4.spv fixtures/turbo4.json
|
||
@echo "[kernels/verify] vulkan turbo3_tcq..."
|
||
./vulkan_verify ../vulkan/turbo3_tcq.spv fixtures/turbo3_tcq.json
|
||
@echo "[kernels/verify] vulkan qjl..."
|
||
./vulkan_verify ../vulkan/qjl.spv fixtures/qjl.json
|
||
@echo "[kernels/verify] vulkan polar..."
|
||
./vulkan_verify ../vulkan/polar.spv fixtures/polar.json
|
||
@echo "[kernels/verify] vulkan polar + QJL residual..."
|
||
./vulkan_verify ../vulkan/polar.spv fixtures/polar_qjl.json
|
||
@echo "[kernels/verify] vulkan polar pre-Hadamard query..."
|
||
./vulkan_verify ../vulkan/polar_preht.spv fixtures/polar.json
|
||
@echo "[kernels/verify] vulkan polar pre-Hadamard query + QJL residual..."
|
||
./vulkan_verify ../vulkan/polar_preht.spv fixtures/polar_qjl.json
|
||
|
||
# Multi-block-per-workgroup Vulkan variants. One SPV family per kernel; the
|
||
# blocks/tokens-per-workgroup specialization constant (constant_id 0) is set
|
||
# at pipeline-create time. Sweeps N ∈ {1,2,4,8}; each must match the same
|
||
# fixture as its base kernel. Additive — not part of the publish gate
|
||
# (vulkan-verify) so a future per-device tuning regression is visible
|
||
# separately from the canonical kernel set.
|
||
vulkan-verify-multiblock: vulkan_verify vulkan-spirv reference-test
|
||
@for n in 1 2 4 8; do \
|
||
echo "[kernels/verify] vulkan turbo3 multi $$n..." && \
|
||
./vulkan_verify ../vulkan/turbo3_multi.spv fixtures/turbo3.json --multi $$n && \
|
||
echo "[kernels/verify] vulkan turbo4 multi $$n..." && \
|
||
./vulkan_verify ../vulkan/turbo4_multi.spv fixtures/turbo4.json --multi $$n && \
|
||
echo "[kernels/verify] vulkan turbo3_tcq multi $$n..." && \
|
||
./vulkan_verify ../vulkan/turbo3_tcq_multi.spv fixtures/turbo3_tcq.json --multi $$n && \
|
||
echo "[kernels/verify] vulkan qjl multi $$n..." && \
|
||
./vulkan_verify ../vulkan/qjl_multi.spv fixtures/qjl.json --multi $$n || exit 1; \
|
||
done
|
||
|
||
# Staged Vulkan fallback entrypoints (decode-to-fp32 / non-attention matvec).
|
||
# Expected outputs are computed in vulkan_verify.cpp from the bit-exact C
|
||
# reference (eliza_qjl_mul_mv / eliza_qjl_dequantize_row / eliza_polar_dequantize_row),
|
||
# so no new fixture file is required. Additive — kept out of the publish gate
|
||
# (vulkan-verify) because these are not on the runtime hot path; the runtime
|
||
# dispatches qjl.comp / polar.comp / polar_preht.comp.
|
||
#
|
||
# * qjl_mul_mv: hard gate. It is the QJL-score parity check (n_heads ==
|
||
# n_kv_heads == 1, n_rows == n_tokens), reduction-based, and passes
|
||
# deterministically on every ICD tested (Intel ANV, lavapipe).
|
||
# * qjl_get_rows / polar_get_rows: informational. These write one fp32
|
||
# output per element from all 32 invocations (no reduction). They are
|
||
# bit-exact and deterministic on lavapipe, but on Mesa ANV (Arrow Lake /
|
||
# Xe-LPG, Mesa 25.2.8) a handful of output elements intermittently read
|
||
# back 0 — an ANV codegen issue with the all-threads strided SSBO store,
|
||
# not a logic bug (the reduction-based polar.comp/polar_preht.comp share
|
||
# the same decode and are clean). Run shown for visibility; `-` so a
|
||
# flaky ANV result does not fail this target.
|
||
vulkan-verify-fallbacks: vulkan_verify vulkan-spirv reference-test
|
||
@echo "[kernels/verify] vulkan qjl_mul_mv (== qjl score parity, n_rows=8)..."
|
||
./vulkan_verify ../vulkan/qjl_mul_mv.spv fixtures/qjl.json
|
||
@echo "[kernels/verify] vulkan qjl_get_rows (decode block 0 vs eliza_qjl_dequantize_row; informational on ANV)..."
|
||
-./vulkan_verify ../vulkan/qjl_get_rows.spv fixtures/qjl.json
|
||
@echo "[kernels/verify] vulkan polar_get_rows use_qjl=0 (decode block 0 vs eliza_polar_dequantize_row; informational on ANV)..."
|
||
-./vulkan_verify ../vulkan/polar_get_rows.spv fixtures/polar.json
|
||
@echo "[kernels/verify] vulkan polar_get_rows use_qjl=1 (informational on ANV)..."
|
||
-./vulkan_verify ../vulkan/polar_get_rows.spv fixtures/polar_qjl.json
|
||
|
||
# Fused-attention (QJL-K + TBQ-V / Polar-V) Vulkan gate. vulkan/fused_attn_qjl_tbq.comp
|
||
# and vulkan/fused_attn_qjl_polar.comp port GGML_OP_FUSED_ATTN_QJL_TBQ + the Polar
|
||
# V-mix variant: one-pass FlashAttention-style online softmax per head, the score
|
||
# vector never materialized, a driver-portable 32-thread shared-memory reduction
|
||
# (no subgroup assumptions). vulkan_verify's run_fused_attn() parses the
|
||
# `cases`-array fixture schema (n_kv ∈ {64, 512, 256, 128} standing in for the
|
||
# 4k/32k/128k/256k regimes, GQA fanouts 1/2/4) and diffs against the C reference
|
||
# (eliza_fused_attn_qjl_tbq3 / eliza_fused_attn_qjl_polar). 1e-3 absolute tolerance.
|
||
# Backend-neutral contract: reports/porting/2026-05-11/fused-attn-op-contract.md
|
||
# Metal-specific design: reports/porting/2026-05-11/metal-fused-attn-and-polar-preht-design.md
|
||
vulkan-verify-fused: vulkan_verify vulkan-spirv reference-test
|
||
@echo "[kernels/verify] vulkan fused-attn QJL-K + TBQ3-V..."
|
||
./vulkan_verify ../vulkan/fused_attn_qjl_tbq.spv fixtures/fused_attn_qjl_tbq.json
|
||
@echo "[kernels/verify] vulkan fused-attn QJL-K + TBQ3-V causal prefix..."
|
||
./vulkan_verify ../vulkan/fused_attn_qjl_tbq.spv fixtures/fused_attn_qjl_tbq_causal.json
|
||
@echo "[kernels/verify] vulkan fused-attn QJL-K + Q4_POLAR-V (use_qjl=1)..."
|
||
./vulkan_verify ../vulkan/fused_attn_qjl_polar.spv fixtures/fused_attn_qjl_polar.json
|
||
@echo "[kernels/verify] vulkan fused-attn QJL-K + Q4_POLAR-V causal prefix (use_qjl=1)..."
|
||
./vulkan_verify ../vulkan/fused_attn_qjl_polar.spv fixtures/fused_attn_qjl_polar_causal.json
|
||
|
||
# Native Linux end-to-end smoke for the Vulkan backend. This first runs the
|
||
# standalone fixture verifier, then builds the patched fork target and runs the
|
||
# graph-dispatch smoke below. It writes hardware-results evidence, rejects
|
||
# software ICDs unless ELIZA_ALLOW_SOFTWARE_VULKAN=1 is set, and refuses stale
|
||
# or symbol-only artifacts.
|
||
vulkan-native-smoke:
|
||
./linux_vulkan_smoke.sh
|
||
|
||
# Android on-device Vulkan smoke. Cross-compiles vulkan_verify with the NDK,
|
||
# pushes the verifier + SPIR-V + fixtures via adb, runs the standalone fixture
|
||
# checks, then runs the built-fork graph-dispatch smoke unless
|
||
# ELIZA_ANDROID_VULKAN_STANDALONE_ONLY=1 is set.
|
||
android-vulkan-smoke:
|
||
./android_vulkan_smoke.sh
|
||
|
||
# Android built-fork graph-dispatch smoke. Builds android-arm64-vulkan through
|
||
# the mtp build script using a standalone cache, cross-compiles
|
||
# vulkan_dispatch_smoke.cpp, pushes it plus libggml-vulkan.so to a physical
|
||
# device, and writes target-specific runtime evidence.
|
||
android-vulkan-graph-smoke:
|
||
./android_vulkan_graph_smoke.sh
|
||
|
||
# Metal harness (macOS only). Requires xcrun on PATH.
|
||
metal: metal_verify
|
||
metal_verify: metal_verify.mm $(REF_OBJ) $(QJL_POLAR_OBJ)
|
||
@which xcrun >/dev/null 2>&1 || { echo "xcrun not found — Metal harness needs macOS Xcode tools."; exit 1; }
|
||
clang++ $(CXXFLAGS) -ObjC++ -fobjc-arc -I../reference \
|
||
metal_verify.mm $(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
-framework Foundation -framework Metal -framework MetalKit -framework CoreGraphics \
|
||
-lm -o metal_verify
|
||
|
||
# Full Metal correctness gate for the default-on kernel set. The single-block
|
||
# kernels are the compatibility entrypoints; Polar is checked both with and
|
||
# without its QJL residual branch.
|
||
metal-verify: metal_verify reference-test
|
||
@echo "[kernels/verify] metal turbo3..."
|
||
./metal_verify ../metal/turbo3.metal kernel_turbo3_dot fixtures/turbo3.json
|
||
@echo "[kernels/verify] metal turbo4..."
|
||
./metal_verify ../metal/turbo4.metal kernel_turbo4_dot fixtures/turbo4.json
|
||
@echo "[kernels/verify] metal turbo3_tcq..."
|
||
./metal_verify ../metal/turbo3_tcq.metal kernel_turbo3_tcq_dot fixtures/turbo3_tcq.json
|
||
@echo "[kernels/verify] metal qjl..."
|
||
./metal_verify ../metal/qjl.metal kernel_attn_score_qjl1_256 fixtures/qjl.json
|
||
@echo "[kernels/verify] metal polar..."
|
||
./metal_verify ../metal/polar.metal kernel_mul_mv_q4_polar_f32 fixtures/polar.json
|
||
@echo "[kernels/verify] metal polar + QJL residual..."
|
||
./metal_verify ../metal/polar.metal kernel_mul_mv_q4_polar_f32 fixtures/polar_qjl.json
|
||
@echo "[kernels/verify] metal polar pre-Hadamard query..."
|
||
./metal_verify ../metal/polar.metal kernel_mul_mv_q4_polar_preht_f32 fixtures/polar.json
|
||
@echo "[kernels/verify] metal polar pre-Hadamard query + QJL residual..."
|
||
./metal_verify ../metal/polar.metal kernel_mul_mv_q4_polar_preht_f32 fixtures/polar_qjl.json
|
||
|
||
# Runtime-source-of-truth gate. metal-verify above tests the verify-harness copy
|
||
# under ../metal/; the SHIPPED kernels the fork actually embeds into
|
||
# default.metallib live in the llama.cpp submodule under eliza-shipped/. Those two
|
||
# copies MUST stay in lockstep — when they drifted (fork commit 412b8487b regressed
|
||
# eliza-shipped/turbo3+turbo4 to 0/8 while ../metal/ stayed 8/8) the metal-verify
|
||
# gate went green but the runtime shipped broken TBQ3_0/TBQ4_0 attention scores.
|
||
# This target runs the same parity checks against the embedded kernels so the
|
||
# regression class fails CI instead of shipping silently.
|
||
SHIPPED_METAL := ../llama.cpp/ggml/src/ggml-metal/eliza-shipped
|
||
metal-verify-shipped: metal_verify reference-test
|
||
@test -d $(SHIPPED_METAL) || { echo "[kernels/verify] $(SHIPPED_METAL) missing — run git submodule update --init"; exit 1; }
|
||
@echo "[kernels/verify] SHIPPED metal turbo3..."
|
||
./metal_verify $(SHIPPED_METAL)/turbo3.metal kernel_turbo3_dot fixtures/turbo3.json
|
||
@echo "[kernels/verify] SHIPPED metal turbo4..."
|
||
./metal_verify $(SHIPPED_METAL)/turbo4.metal kernel_turbo4_dot fixtures/turbo4.json
|
||
@echo "[kernels/verify] SHIPPED metal turbo3_tcq..."
|
||
./metal_verify $(SHIPPED_METAL)/turbo3_tcq.metal kernel_turbo3_tcq_dot fixtures/turbo3_tcq.json
|
||
@echo "[kernels/verify] SHIPPED metal qjl..."
|
||
./metal_verify $(SHIPPED_METAL)/qjl.metal kernel_attn_score_qjl1_256 fixtures/qjl.json
|
||
@echo "[kernels/verify] SHIPPED metal polar (+QJL residual)..."
|
||
./metal_verify $(SHIPPED_METAL)/polar.metal kernel_mul_mv_q4_polar_f32 fixtures/polar.json
|
||
./metal_verify $(SHIPPED_METAL)/polar.metal kernel_mul_mv_q4_polar_f32 fixtures/polar_qjl.json
|
||
@echo "[kernels/verify] SHIPPED metal fused QJL-K + TBQ3-V..."
|
||
./metal_verify $(SHIPPED_METAL)/fused_attn_qjl_tbq.metal kernel_fused_attn_qjl_tbq3_f32 fixtures/fused_attn_qjl_tbq.json
|
||
@echo "[kernels/verify] SHIPPED metal fused QJL-K + Q4_POLAR-V..."
|
||
./metal_verify $(SHIPPED_METAL)/fused_attn_qjl_polar.metal kernel_fused_attn_qjl_polar_f32 fixtures/fused_attn_qjl_polar.json
|
||
@echo "[kernels/verify] SHIPPED metal set PASS — runtime-embedded kernels match the C reference"
|
||
|
||
# Additive check for the long-context multi-block Metal entrypoints. These
|
||
# are separate symbols from the compatibility entrypoints and need direct
|
||
# harness coverage whenever their dispatch shape changes.
|
||
metal-verify-multiblock: metal_verify reference-test
|
||
@for n in 2 3 4 8; do \
|
||
echo "[kernels/verify] metal turbo3 multi $$n..." && \
|
||
./metal_verify ../metal/turbo3.metal kernel_turbo3_dot_multi fixtures/turbo3.json --multi $$n && \
|
||
echo "[kernels/verify] metal turbo4 multi $$n..." && \
|
||
./metal_verify ../metal/turbo4.metal kernel_turbo4_dot_multi fixtures/turbo4.json --multi $$n && \
|
||
echo "[kernels/verify] metal turbo3_tcq multi $$n..." && \
|
||
./metal_verify ../metal/turbo3_tcq.metal kernel_turbo3_tcq_dot_multi fixtures/turbo3_tcq.json --multi $$n && \
|
||
echo "[kernels/verify] metal qjl multi $$n..." && \
|
||
./metal_verify ../metal/qjl.metal kernel_attn_score_qjl1_256_multi fixtures/qjl.json --multi $$n || exit 1; \
|
||
done
|
||
|
||
# Fused-attention (QJL-K + TBQ-V / Polar-V) + Polar pre-Hadamard-query Metal
|
||
# gate. metal/fused_attn_qjl_tbq.metal and metal/fused_attn_qjl_polar.metal port
|
||
# GGML_OP_FUSED_ATTN_QJL_TBQ + the Polar V-mix variant (one-pass online softmax,
|
||
# per-token QJL score via simd_sum, threadgroup-scratch TBQ Hadamard-32 / Polar
|
||
# Hadamard-128, score vector never materialized). metal/polar_preht.metal adds
|
||
# kernel_attn_score_q4_polar_preht_f32(_multi) — Polar dot vs a pre-Hadamarded
|
||
# query. metal_verify's run_fused_attn() / run_polar_preht() parse the cases-array
|
||
# and the polar_preht fixture; both diff bit-for-bit against the C references
|
||
# (eliza_fused_attn_qjl_tbq3 / eliza_fused_attn_qjl_polar / eliza_polar_mul_mv).
|
||
# 1e-3 absolute tolerance.
|
||
#
|
||
# Hardware-verified on Apple M4 Max: the kernels are byte-faithful mirrors of the
|
||
# hardware-verified Vulkan ports (vulkan/fused_attn_*.comp, vulkan/polar_preht.comp,
|
||
# Intel ARL Mesa ANV). This standalone fixture gate still does not make the
|
||
# fused_attn Metal capability runtime-ready; that requires a built-fork Metal
|
||
# graph-dispatch smoke. fused_attn is an optimization on top of the five required
|
||
# kernels — AGENTS.md §3 — not a required kernel.
|
||
# Backend-neutral contract: reports/porting/2026-05-11/fused-attn-op-contract.md
|
||
# Metal-specific design: reports/porting/2026-05-11/metal-fused-attn-and-polar-preht-design.md
|
||
metal-verify-fused: metal_verify reference-test
|
||
@echo "[kernels/verify] metal fused-attn QJL-K + TBQ3-V..."
|
||
./metal_verify ../metal/fused_attn_qjl_tbq.metal kernel_fused_attn_qjl_tbq3_f32 fixtures/fused_attn_qjl_tbq.json
|
||
@echo "[kernels/verify] metal fused-attn QJL-K + TBQ3-V causal prefix..."
|
||
./metal_verify ../metal/fused_attn_qjl_tbq.metal kernel_fused_attn_qjl_tbq3_f32 fixtures/fused_attn_qjl_tbq_causal.json
|
||
@echo "[kernels/verify] metal fused-attn QJL-K + Q4_POLAR-V (use_qjl=1)..."
|
||
./metal_verify ../metal/fused_attn_qjl_polar.metal kernel_fused_attn_qjl_polar_f32 fixtures/fused_attn_qjl_polar.json
|
||
@echo "[kernels/verify] metal fused-attn QJL-K + Q4_POLAR-V causal prefix (use_qjl=1)..."
|
||
./metal_verify ../metal/fused_attn_qjl_polar.metal kernel_fused_attn_qjl_polar_f32 fixtures/fused_attn_qjl_polar_causal.json
|
||
@echo "[kernels/verify] metal Polar pre-Hadamard query score (use_qjl 0/1)..."
|
||
./metal_verify ../metal/polar_preht.metal kernel_attn_score_q4_polar_preht_f32 fixtures/polar_preht.json
|
||
@for n in 2 3 4 8; do \
|
||
echo "[kernels/verify] metal Polar pre-Hadamard query score multi $$n..." && \
|
||
./metal_verify ../metal/polar_preht.metal kernel_attn_score_q4_polar_preht_f32_multi fixtures/polar_preht.json --multi $$n || exit 1; \
|
||
done
|
||
|
||
# Metal performance harness (macOS only). Sibling of metal_verify; measures
|
||
# end-to-end GPU latency, CPU wait, and bandwidth utilisation per kernel at
|
||
# production-relevant sizes. Does NOT replace metal_verify — both build
|
||
# independently against the same shader files.
|
||
metal-bench: metal_bench
|
||
metal_bench: metal_bench.mm
|
||
@which xcrun >/dev/null 2>&1 || { echo "xcrun not found — Metal bench needs macOS Xcode tools."; exit 1; }
|
||
clang++ $(CXXFLAGS) -ObjC++ -fobjc-arc \
|
||
metal_bench.mm \
|
||
-framework Foundation -framework Metal -framework MetalKit -framework CoreGraphics \
|
||
-lm -o metal_bench
|
||
|
||
# Convenience target: build metal_bench then run --mode batched. The batched
|
||
# sweep dispatches N ∈ {1,4,16,64,128,256} into a single MTLCommandBuffer
|
||
# per kernel and reports per-dispatch GPU cost vs N. See
|
||
# bench_M4Max_batched_2026-05-10.md for the Wave-6 finding.
|
||
metal-bench-batched: metal_bench
|
||
@mkdir -p bench_results
|
||
./metal_bench --mode batched --out bench_results/m4max_batched_2026-05-10.json
|
||
|
||
# Multi-block-per-dispatch sweep (SHADER_REVIEW M3). Each threadgroup processes
|
||
# N consecutive KV blocks serially via 32-thread loop; cuts launch grid by N×.
|
||
# Sweeps the 4 small kernels (turbo3, turbo4, turbo3_tcq, qjl) over
|
||
# N ∈ {1,2,4,8,16,32} and writes per-kernel optimal-N + speedup.
|
||
metal-bench-multiblock: metal_bench
|
||
@mkdir -p bench_results
|
||
./metal_bench --mode multiblock --out bench_results/m4max_multiblock_2026-05-10.json
|
||
|
||
# Fused-attention benchmark for the standalone Metal QJL-K + quantized-V
|
||
# kernels. This complements metal-verify-fused: correctness stays in
|
||
# metal_verify, while this target records the one-pass online-softmax kernel
|
||
# cost at a 9B-class single-token decode shape.
|
||
metal-bench-fused: metal_bench
|
||
@mkdir -p bench_results
|
||
./metal_bench --mode fused --out bench_results/m4max_fused_2026-05-12.json
|
||
|
||
# CPU reference baseline. Single-threaded scalar C, no GPU. Same workload
|
||
# constants as metal_bench so the BENCHMARK doc can join the two tables.
|
||
HOST_UNAME_S := $(shell uname -s)
|
||
HOST_UNAME_M := $(shell uname -m)
|
||
HOST_MTP_TARGET := $(if $(filter Darwin,$(HOST_UNAME_S)),darwin-arm64-metal-fused,$(if $(filter aarch64 arm64,$(HOST_UNAME_M)),linux-aarch64-cpu,linux-x64-cpu))
|
||
ELIZA_MTP_LLAMA_DIR ?= ../llama.cpp
|
||
ELIZA_MTP_CPU_BIN_DIR ?= $(firstword \
|
||
$(wildcard $(ELIZA_MTP_LLAMA_DIR)/build/$(HOST_MTP_TARGET)/bin) \
|
||
$(wildcard $(HOME)/.eliza/local-inference/bin/mtp/$(HOST_MTP_TARGET)) \
|
||
$(wildcard $(ELIZA_MTP_LLAMA_DIR)/build/linux-x64-cpu/bin) \
|
||
$(wildcard $(HOME)/.eliza/local-inference/bin/mtp/linux-x64-cpu))
|
||
GGML_INC_DIR ?= $(ELIZA_MTP_LLAMA_DIR)/ggml/include
|
||
CPU_DISPATCH_SMOKE_BIN ?= qjl_mt_check.$(HOST_MTP_TARGET)
|
||
|
||
$(CPU_DISPATCH_SMOKE_BIN): qjl_mt_check.c
|
||
@test -n "$(ELIZA_MTP_CPU_BIN_DIR)" || { echo "cpu-dispatch-smoke needs a built llama.cpp CPU-compatible binary dir (set ELIZA_MTP_CPU_BIN_DIR)."; exit 1; }
|
||
$(CC) $(CFLAGS) -I$(GGML_INC_DIR) qjl_mt_check.c \
|
||
-L$(ELIZA_MTP_CPU_BIN_DIR) -lggml -lggml-base -lggml-cpu \
|
||
-Wl,-rpath,$(ELIZA_MTP_CPU_BIN_DIR) -lm -o $(CPU_DISPATCH_SMOKE_BIN)
|
||
|
||
cpu-dispatch-smoke: $(CPU_DISPATCH_SMOKE_BIN)
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_CPU_BIN_DIR):$${DYLD_LIBRARY_PATH:-} \
|
||
LD_LIBRARY_PATH=$(ELIZA_MTP_CPU_BIN_DIR):$${LD_LIBRARY_PATH:-} \
|
||
./$(CPU_DISPATCH_SMOKE_BIN)
|
||
|
||
# QJL1_256-K / Q4_POLAR-V flash_attn_ext dispatch smoke. Builds the post-cast
|
||
# (F16 K/V) attention graph through ggml_flash_attn_ext and asserts finite,
|
||
# non-zero output — guards the dequant-hop path in src/llama-graph.cpp from
|
||
# regressing back to a NULL-vec_dot call on the K-cache-only quant types.
|
||
# Needs a built ggml CPU lib (set ELIZA_MTP_CPU_BIN_DIR + GGML_INC_DIR; e.g.
|
||
# a stock llama.cpp build: ELIZA_MTP_CPU_BIN_DIR=<build>/bin
|
||
# GGML_INC_DIR=<src>/ggml/include).
|
||
CPU_QJL_POLAR_ATTN_SMOKE_BIN ?= cpu_qjl_polar_attn_smoke.$(HOST_MTP_TARGET)
|
||
|
||
$(CPU_QJL_POLAR_ATTN_SMOKE_BIN): cpu_qjl_polar_attn_smoke.c
|
||
@test -n "$(ELIZA_MTP_CPU_BIN_DIR)" || { echo "cpu-qjl-polar-attn-smoke needs a built llama.cpp CPU-compatible binary dir (set ELIZA_MTP_CPU_BIN_DIR)."; exit 1; }
|
||
$(CC) $(CFLAGS) -I$(GGML_INC_DIR) cpu_qjl_polar_attn_smoke.c \
|
||
-L$(ELIZA_MTP_CPU_BIN_DIR) -lggml -lggml-base -lggml-cpu \
|
||
-Wl,-rpath,$(ELIZA_MTP_CPU_BIN_DIR) -lm -o $(CPU_QJL_POLAR_ATTN_SMOKE_BIN)
|
||
|
||
cpu-qjl-polar-attn-smoke: $(CPU_QJL_POLAR_ATTN_SMOKE_BIN)
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_CPU_BIN_DIR):$${DYLD_LIBRARY_PATH:-} \
|
||
LD_LIBRARY_PATH=$(ELIZA_MTP_CPU_BIN_DIR):$${LD_LIBRARY_PATH:-} \
|
||
./$(CPU_QJL_POLAR_ATTN_SMOKE_BIN)
|
||
|
||
cpu-bench: cpu_bench
|
||
cpu_bench: cpu_bench.c $(REF_OBJ) $(QJL_POLAR_OBJ)
|
||
$(CC) $(CFLAGS) -I../reference cpu_bench.c $(REF_OBJ) $(QJL_POLAR_OBJ) -lm -o cpu_bench
|
||
|
||
# CPU SIMD-path bench: links the qjl-cpu / polarquant-cpu plugin static
|
||
# libraries and times their dispatched SIMD entrypoints (AVX-VNNI / AVX2
|
||
# / NEON-dotprod) at the production attention workload, with a thread
|
||
# sweep. Requires the plugin libs — built here via cmake first.
|
||
QJL_CPU_DIR = ../../../../packages/native-plugins/qjl-cpu
|
||
POLAR_CPU_DIR = ../../../../packages/native-plugins/polarquant-cpu
|
||
QJL_CPU_LIB = $(QJL_CPU_DIR)/build/libqjl.a
|
||
POLAR_CPU_LIB = $(POLAR_CPU_DIR)/build/libpolarquant.a
|
||
|
||
$(QJL_CPU_LIB):
|
||
cmake -S $(QJL_CPU_DIR) -B $(QJL_CPU_DIR)/build -DCMAKE_BUILD_TYPE=Release >/dev/null
|
||
cmake --build $(QJL_CPU_DIR)/build --target qjl >/dev/null
|
||
$(POLAR_CPU_LIB):
|
||
cmake -S $(POLAR_CPU_DIR) -B $(POLAR_CPU_DIR)/build -DCMAKE_BUILD_TYPE=Release >/dev/null
|
||
cmake --build $(POLAR_CPU_DIR)/build --target polarquant >/dev/null
|
||
|
||
cpu-simd-bench: cpu_simd_bench
|
||
OPENMP_FLAGS ?= $(shell printf '' | $(CC) -x c -fopenmp -E - >/dev/null 2>&1 && printf '%s' '-fopenmp')
|
||
THREAD_FLAGS ?= $(if $(OPENMP_FLAGS),$(OPENMP_FLAGS),-pthread)
|
||
|
||
cpu_simd_bench: cpu_simd_bench.c $(QJL_CPU_LIB) $(POLAR_CPU_LIB)
|
||
$(CC) -O3 -march=native $(THREAD_FLAGS) -std=c11 \
|
||
-I$(QJL_CPU_DIR)/include -I$(QJL_CPU_DIR)/src -I$(POLAR_CPU_DIR)/include \
|
||
cpu_simd_bench.c $(QJL_CPU_LIB) $(POLAR_CPU_LIB) -lm -o cpu_simd_bench
|
||
|
||
# Umbrella: build (and run, into bench_results/) every available bench.
|
||
# CPU bench always runs. Metal bench runs only on macOS with xcrun.
|
||
bench: cpu-bench
|
||
@mkdir -p bench_results
|
||
@echo "[bench] running cpu_bench ..."
|
||
./cpu_bench --runs 3 --warmup 1 --out bench_results/cpu_m4max_2026-05-10.json
|
||
@if which xcrun >/dev/null 2>&1; then \
|
||
echo "[bench] running metal_bench ..." ; \
|
||
$(MAKE) metal-bench >/dev/null && \
|
||
./metal_bench --out bench_results/m4max_2026-05-10.json && \
|
||
echo "[bench] running metal_bench --mode tgsweep ..." && \
|
||
./metal_bench --mode tgsweep --out bench_results/m4max_tgsweep_2026-05-10.json && \
|
||
echo "[bench] running metal_bench --mode fp16ref ..." && \
|
||
./metal_bench --mode fp16ref --out bench_results/m4max_fp16ref_2026-05-10.json && \
|
||
echo "[bench] running metal_bench --mode multiblock ..." && \
|
||
./metal_bench --mode multiblock --out bench_results/m4max_multiblock_2026-05-10.json && \
|
||
echo "[bench] running metal_bench --mode fused ..." && \
|
||
./metal_bench --mode fused --out bench_results/m4max_fused_2026-05-12.json ; \
|
||
else \
|
||
echo "[bench] skipping metal_bench (xcrun not found — macOS only)" ; \
|
||
fi
|
||
|
||
# Compile the GLSL compute shaders to SPIR-V. The vulkan_verify harness loads
|
||
# .spv files. glslc resolution: prefer system glslc, fall back to the Android
|
||
# NDK's bundled glslc (same path the mtp build script uses).
|
||
GLSLC ?= $(shell command -v glslc 2>/dev/null)
|
||
ifeq ($(GLSLC),)
|
||
NDK_GLSLC := $(firstword \
|
||
$(wildcard $(HOME)/Library/Android/sdk/ndk/*/shader-tools/darwin-x86_64/glslc) \
|
||
$(wildcard $(HOME)/Library/Android/sdk/ndk/*/shader-tools/linux-x86_64/glslc) \
|
||
$(wildcard $(HOME)/Android/Sdk/ndk/*/shader-tools/linux-x86_64/glslc) \
|
||
$(wildcard $(HOME)/Android/Sdk/ndk/*/shader-tools/darwin-x86_64/glslc))
|
||
GLSLC := $(NDK_GLSLC)
|
||
endif
|
||
|
||
VULKAN_SHADERS = \
|
||
../vulkan/turbo3.spv \
|
||
../vulkan/turbo4.spv \
|
||
../vulkan/turbo3_tcq.spv \
|
||
../vulkan/qjl.spv \
|
||
../vulkan/qjl_get_rows.spv \
|
||
../vulkan/qjl_mul_mv.spv \
|
||
../vulkan/polar.spv \
|
||
../vulkan/polar_preht.spv \
|
||
../vulkan/polar_get_rows.spv \
|
||
../vulkan/turbo3_multi.spv \
|
||
../vulkan/turbo4_multi.spv \
|
||
../vulkan/turbo3_tcq_multi.spv \
|
||
../vulkan/qjl_multi.spv \
|
||
../vulkan/fused_attn_qjl_tbq.spv \
|
||
../vulkan/fused_attn_qjl_polar.spv \
|
||
../vulkan/mul_mm_scalar.spv \
|
||
../vulkan/mul_mm_tiled.spv \
|
||
../vulkan/mul_mm_reg.spv
|
||
|
||
vulkan-spirv: $(VULKAN_SHADERS)
|
||
|
||
../vulkan/%.spv: ../vulkan/%.comp
|
||
ifeq ($(GLSLC),)
|
||
@echo "glslc not found — install LunarG SDK or Android NDK to build SPIR-V."
|
||
@false
|
||
else
|
||
# subgroupAdd needs SPIR-V 1.3 + Vulkan 1.1 (the kernels target Vulkan 1.1
|
||
# subgroups via GL_KHR_shader_subgroup_arithmetic).
|
||
$(GLSLC) --target-env=vulkan1.1 --target-spv=spv1.3 \
|
||
-fshader-stage=compute $< -o $@
|
||
endif
|
||
|
||
# CUDA harness. Sibling of metal/vulkan_verify. Self-contained: needs only
|
||
# nvcc + the CUDA runtime headers (the in-fork ggml-cuda headers are NOT a
|
||
# build dependency — the harness ports the same reference algorithm into CUDA
|
||
# __device__ kernels that consume the fixture byte images directly, exactly
|
||
# like metal_verify / vulkan_verify port it into the shader). It links the
|
||
# same qjl_polar_ref.o the Metal/Vulkan harnesses link, for an on-host
|
||
# cross-check of the fixture's expected_scores.
|
||
#
|
||
# Resolution order:
|
||
# 1. nvcc on PATH — required to compile (.cu -> ELF).
|
||
# 2. CUDA_HOME/include — defaults to /usr/local/cuda; overridable.
|
||
# 3. ELIZA_MTP_LLAMA_DIR / ELIZA_MTP_LIBGGML_CUDA — UNUSED by this
|
||
# harness now; retained as env so cuda_runner.sh's passthrough is inert.
|
||
NVCC ?= $(shell command -v nvcc 2>/dev/null)
|
||
# Distro CUDA toolkits (Debian/Ubuntu `nvidia-cuda-toolkit`) put the runtime
|
||
# headers in /usr/include, not /usr/local/cuda/include. Probe both.
|
||
CUDA_HOME ?= $(if $(wildcard /usr/local/cuda/include/cuda_runtime.h),/usr/local/cuda,$(if $(wildcard /usr/include/cuda_runtime.h),/usr,/usr/local/cuda))
|
||
ELIZA_MTP_IN_TREE_LLAMA_DIR := $(abspath ../llama.cpp)
|
||
# Prefer the checked-out submodule when running from this repository. The
|
||
# ~/.cache build clone can be stale relative to the active worktree, and graph
|
||
# dispatch smokes must link against the headers/libs matching the source under
|
||
# review. CI/external runners can still override ELIZA_MTP_LLAMA_DIR.
|
||
ELIZA_MTP_LLAMA_DIR ?= $(if $(wildcard $(ELIZA_MTP_IN_TREE_LLAMA_DIR)/ggml/include/ggml.h),$(ELIZA_MTP_IN_TREE_LLAMA_DIR),$(HOME)/.cache/eliza-mtp/eliza-llama-cpp)
|
||
ELIZA_MTP_LIBGGML_CUDA ?= $(ELIZA_MTP_LLAMA_DIR)/build-cuda/ggml/src/ggml-cuda/libggml-cuda.so
|
||
|
||
# CMAKE_CUDA_ARCHITECTURES the fork build hook should pin (platform-targets
|
||
# agent owns build-llama-cpp-mtp.mjs):
|
||
# linux-x64-cuda : 80;86;89;90;100;120 (Ampere..Blackwell, mobile dGPU = 12.0)
|
||
# linux-aarch64-cuda : 90a;90;100 (GH200 sm_90a + Hopper/Blackwell)
|
||
# For a one-off local build on this machine's Blackwell mobile dGPU:
|
||
# ELIZA_MTP_CMAKE_FLAGS='-DCMAKE_CUDA_ARCHITECTURES=120;90' \
|
||
# node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target linux-x64-cuda
|
||
#
|
||
# Compile-time arch list for the harness itself (cheap; one TU). Overridable
|
||
# with CUDA_ARCH_FLAGS. Always emits sm_80/86/90 SASS plus compute_90 PTX —
|
||
# the driver JITs the compute_90 PTX for any newer arch (Blackwell sm_120 on
|
||
# RTX 50xx), so the harness runs on every NVIDIA part since Ampere even with a
|
||
# pre-12.8 nvcc that cannot target sm_100/sm_120 directly. nvcc >= 12.8 adds
|
||
# native sm_120 SASS on top (faster JIT-free first launch on Blackwell consumer).
|
||
comma := ,
|
||
CUDA_SM89_GENCODE := $(shell echo 'int main(){return 0;}' > /tmp/.nvcc_sm89_probe.cu 2>/dev/null && $(NVCC) -arch=sm_89 -c /tmp/.nvcc_sm89_probe.cu -o /dev/null >/dev/null 2>&1 && echo "-gencode arch=compute_89@code=sm_89"; rm -f /tmp/.nvcc_sm89_probe.cu)
|
||
CUDA_SM90A_GENCODE := $(shell echo 'int main(){return 0;}' > /tmp/.nvcc_sm90a_probe.cu 2>/dev/null && $(NVCC) -arch=sm_90a -c /tmp/.nvcc_sm90a_probe.cu -o /dev/null >/dev/null 2>&1 && echo "-gencode arch=compute_90a@code=sm_90a"; rm -f /tmp/.nvcc_sm90a_probe.cu)
|
||
CUDA_SM100_GENCODE := $(shell echo 'int main(){return 0;}' > /tmp/.nvcc_sm100_probe.cu 2>/dev/null && $(NVCC) -arch=sm_100 -c /tmp/.nvcc_sm100_probe.cu -o /dev/null >/dev/null 2>&1 && echo "-gencode arch=compute_100@code=sm_100"; rm -f /tmp/.nvcc_sm100_probe.cu)
|
||
CUDA_BLACKWELL_GENCODE := $(shell echo 'int main(){return 0;}' > /tmp/.nvcc_sm120_probe.cu 2>/dev/null && $(NVCC) -arch=sm_120 -c /tmp/.nvcc_sm120_probe.cu -o /dev/null >/dev/null 2>&1 && echo "-gencode arch=compute_120@code=sm_120"; rm -f /tmp/.nvcc_sm120_probe.cu)
|
||
CUDA_ARCH_FLAGS ?= -gencode arch=compute_80,code=sm_80 \
|
||
-gencode arch=compute_86,code=sm_86 \
|
||
-gencode arch=compute_90,code=sm_90 \
|
||
-gencode arch=compute_90,code=compute_90 \
|
||
$(subst @,$(comma),$(CUDA_SM89_GENCODE)) \
|
||
$(subst @,$(comma),$(CUDA_SM90A_GENCODE)) \
|
||
$(subst @,$(comma),$(CUDA_SM100_GENCODE)) \
|
||
$(subst @,$(comma),$(CUDA_BLACKWELL_GENCODE))
|
||
|
||
CUDA_INCLUDES = -I$(CUDA_HOME)/include -I.
|
||
|
||
cuda: cuda_verify
|
||
cuda_verify: cuda_verify.cu tbq3_tcq_codebook.inc $(QJL_POLAR_OBJ) $(REF_OBJ)
|
||
ifeq ($(NVCC),)
|
||
@echo "CUDA toolchain not found — install via 'apt install nvidia-cuda-toolkit' (Linux) or download the CUDA Toolkit (macOS not supported)."
|
||
@false
|
||
else
|
||
$(NVCC) -O2 -std=c++17 $(CUDA_ARCH_FLAGS) \
|
||
$(CUDA_INCLUDES) \
|
||
cuda_verify.cu $(QJL_POLAR_OBJ) $(REF_OBJ) \
|
||
-lcudart -lm -o cuda_verify
|
||
endif
|
||
|
||
# Source-level correctness check (a non-CUDA host can run this — nvcc NOT
|
||
# required). Two phases:
|
||
# 1. Harness self-consistency — assert cuda_verify.cu defines a CUDA kernel
|
||
# for every fixture family and includes the generated TCQ codebook.
|
||
# 2. Codebook integrity — assert the generated .inc has exactly 512 floats
|
||
# matching reference/turbo_kernels.c.
|
||
# Full nvcc compile + runtime 8/8 PASS are gated on a CUDA host — see
|
||
# CUDA_VERIFICATION.md.
|
||
cuda-preprocess-check:
|
||
@echo "[cuda-preprocess-check] (1) harness kernel coverage"
|
||
@grep -q 'turbo34_score_kernel' cuda_verify.cu || { echo " MISS: turbo34_score_kernel"; exit 1; }
|
||
@grep -q 'turbo3_tcq_score_kernel' cuda_verify.cu || { echo " MISS: turbo3_tcq_score_kernel"; exit 1; }
|
||
@grep -q 'qjl_score_kernel' cuda_verify.cu || { echo " MISS: qjl_score_kernel"; exit 1; }
|
||
@grep -q 'polar_score_kernel' cuda_verify.cu || { echo " MISS: polar_score_kernel"; exit 1; }
|
||
@grep -q 'tbq3_tcq_codebook.inc' cuda_verify.cu || { echo " MISS: codebook include"; exit 1; }
|
||
@echo " OK: turbo3/turbo4/turbo3_tcq/qjl/polar kernels + codebook include present"
|
||
@echo "[cuda-preprocess-check] (2) generated codebook integrity"
|
||
@test -f tbq3_tcq_codebook.inc || { echo " MISS: tbq3_tcq_codebook.inc — run 'make tbq3_tcq_codebook.inc'"; exit 1; }
|
||
@n=$$(grep -oE '[-+]?[0-9][0-9.eE+-]*f?' tbq3_tcq_codebook.inc | wc -l); \
|
||
test "$$n" = "512" || { echo " FAIL: codebook has $$n entries, expected 512"; exit 1; }
|
||
@echo " OK: tbq3_tcq_codebook.inc has 512 entries"
|
||
@echo "[cuda-preprocess-check] (3) fork CUDA kernel headers present (informational)"
|
||
@for h in turboquant.cuh turbo-tcq.cuh qjl.cuh polarquant.cuh; do \
|
||
if [ -f "$(ELIZA_MTP_LLAMA_DIR)/ggml/src/ggml-cuda/$$h" ]; then echo " present: $$h"; \
|
||
else echo " ABSENT: $$h (fork CUDA KV-cache kernels not yet landed — see CUDA_VERIFICATION.md)"; fi; \
|
||
done
|
||
@echo "[cuda-preprocess-check] OK"
|
||
|
||
# Regenerate the device-side TCQ codebook include from the C reference (the
|
||
# single source of truth: reference/turbo_kernels.c::ELIZA_TURBO3_TCQ_CODEBOOK).
|
||
tbq3_tcq_codebook.inc: ../reference/turbo_kernels.c
|
||
@python3 -c "import re,sys; \
|
||
src=open('../reference/turbo_kernels.c').read(); \
|
||
m=re.search(r'ELIZA_TURBO3_TCQ_CODEBOOK\[512\]\s*=\s*\{(.*?)\};', src, re.S); \
|
||
open('tbq3_tcq_codebook.inc','w').write(m.group(1).strip()+'\n'); \
|
||
n=len(re.findall(r'[-+]?[0-9][0-9.eE+-]*f?', m.group(1))); \
|
||
sys.exit(0 if n==512 else 'codebook extraction got %d entries'%n)"
|
||
@echo "[tbq3_tcq_codebook.inc] regenerated from reference/turbo_kernels.c (512 entries)"
|
||
|
||
# `cuda-verify` runs the harness against all six fixtures, same shape as
|
||
# `bench` does for the perf harness. Gated on nvcc presence so it stays
|
||
# safe on hosts without a CUDA toolchain. (fused_attn_qjl_tbq.json is run too
|
||
# if the fused-attention-reference agent has landed it.)
|
||
cuda-verify:
|
||
ifeq ($(NVCC),)
|
||
@echo "CUDA toolchain not found — install via 'apt install nvidia-cuda-toolkit' (Linux) or download the CUDA Toolkit (macOS not supported)."
|
||
@false
|
||
else
|
||
$(MAKE) cuda
|
||
@echo "[cuda-verify] turbo3 ..." ; ./cuda_verify fixtures/turbo3.json
|
||
@echo "[cuda-verify] turbo4 ..." ; ./cuda_verify fixtures/turbo4.json
|
||
@echo "[cuda-verify] turbo3_tcq ..." ; ./cuda_verify fixtures/turbo3_tcq.json
|
||
@echo "[cuda-verify] qjl ..." ; ./cuda_verify fixtures/qjl.json
|
||
@echo "[cuda-verify] polar ..." ; ./cuda_verify fixtures/polar.json
|
||
@echo "[cuda-verify] polar + QJL residual ..." ; ./cuda_verify fixtures/polar_qjl.json
|
||
@if [ -f fixtures/fused_attn_qjl_tbq.json ]; then \
|
||
echo "[cuda-verify] fused QJL-K + TBQ-V attention ..." ; ./cuda_verify fixtures/fused_attn_qjl_tbq.json ; \
|
||
else echo "[cuda-verify] fused_attn_qjl_tbq.json not present — fused-attention-reference agent owns it; skipping" ; fi
|
||
endif
|
||
|
||
# Focused fused-attention CUDA gate. cuda_verify links a fixture-parity harness
|
||
# with a fused_attn_qjl_tbq3_kernel (warp reductions, online softmax, TCQ
|
||
# codebook in __constant__, page tiling) + the DP4A int8-QJL path; the
|
||
# production shape lives in cuda/fused-attn-qjl-tbq.cu (AUTHORED — hardware-verify
|
||
# pending; mirror it into the fork's ggml-cuda/ via the kernel-patches hook and
|
||
# build with -DGGML_CUDA_FUSED_ATTN_QJL=ON). Run on an NVIDIA host via
|
||
# `verify/cuda_runner.sh` — no NVIDIA hardware on the authoring machine.
|
||
# fused_attn stays out of requiredRuntimeCapabilityKeys (AGENTS.md §3).
|
||
cuda-verify-fused:
|
||
ifeq ($(NVCC),)
|
||
@echo "CUDA toolchain not found — run verify/cuda_runner.sh on an NVIDIA host."
|
||
@false
|
||
else
|
||
$(MAKE) cuda
|
||
@if [ -f fixtures/fused_attn_qjl_tbq.json ]; then \
|
||
echo "[cuda-verify-fused] fused QJL-K + TBQ-V attention (+ DP4A QJL cross-check) ..." ; ./cuda_verify fixtures/fused_attn_qjl_tbq.json ; \
|
||
else echo "[cuda-verify-fused] fused_attn_qjl_tbq.json missing — run reference-test first" ; false ; fi
|
||
endif
|
||
|
||
# Hardware runners. These intentionally fail without matching physical
|
||
# hardware and a model-backed graph smoke (ELIZA_MTP_SMOKE_MODEL).
|
||
cuda-hardware:
|
||
./cuda_runner.sh
|
||
|
||
rocm-hardware:
|
||
./rocm_runner.sh
|
||
|
||
gh200-hardware:
|
||
./gh200_runner.sh
|
||
|
||
windows-hardware:
|
||
@echo "Run on a native Windows host:"
|
||
@echo " pwsh -File packages/inference/verify/windows_runner.ps1 -Backend cuda -Model C:\\\\models\\\\eliza-1-smoke.gguf"
|
||
|
||
# Drift check between the standalone shaders in packages/inference/{metal,vulkan}/
|
||
# and the fork's in-tree shaders. Both ship the same kernel symbols and decode
|
||
# math (the standalones add perf-only divergence; see PATCH_AUDIT_2026-05-10.md
|
||
# §"Standalone-vs-fork shader drift"). The reference runs of metal_verify /
|
||
# vulkan_verify validated the standalones; this target re-runs the same harness
|
||
# against the fork's in-tree variants so we can confirm 8/8 PASS still holds
|
||
# before signing off a bundle.
|
||
#
|
||
# The fork is the elizaOS/llama.cpp checkout the build script clones into
|
||
# ~/.cache/eliza-mtp/eliza-llama-cpp. If the cache dir is absent, this
|
||
# target prints a clear "fork not present, please run build first" message
|
||
# rather than failing cryptically.
|
||
ELIZA_MTP_FORK_METAL_DIR := $(ELIZA_MTP_LLAMA_DIR)/ggml/src/ggml-metal/eliza-shipped
|
||
ELIZA_MTP_FORK_VULKAN_DIR := $(ELIZA_MTP_LLAMA_DIR)/ggml/src/ggml-vulkan/vulkan-shaders
|
||
|
||
verify-fork:
|
||
ifeq ($(wildcard $(ELIZA_MTP_LLAMA_DIR)/.git),)
|
||
@echo "[verify-fork] fork not present at $(ELIZA_MTP_LLAMA_DIR)"
|
||
@echo "[verify-fork] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target darwin-arm64-metal'"
|
||
@echo "[verify-fork] (or any other target) first to clone it, then re-run this target."
|
||
@false
|
||
else
|
||
@if which xcrun >/dev/null 2>&1; then \
|
||
if [ ! -d "$(ELIZA_MTP_FORK_METAL_DIR)" ]; then \
|
||
echo "[verify-fork] fork present but $(ELIZA_MTP_FORK_METAL_DIR) missing — fork layout drift."; \
|
||
echo "[verify-fork] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target darwin-arm64-metal --dry-run' to stage/describe the shipped kernel patch."; \
|
||
exit 1; \
|
||
fi; \
|
||
$(MAKE) metal_verify >/dev/null && \
|
||
echo "[verify-fork] metal/turbo3 (patched shipped: turbo3.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/turbo3.metal kernel_turbo3_dot fixtures/turbo3.json && \
|
||
echo "[verify-fork] metal/turbo4 (patched shipped: turbo4.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/turbo4.metal kernel_turbo4_dot fixtures/turbo4.json && \
|
||
echo "[verify-fork] metal/turbo3_tcq (patched shipped: turbo3_tcq.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/turbo3_tcq.metal kernel_turbo3_tcq_dot fixtures/turbo3_tcq.json && \
|
||
echo "[verify-fork] metal/qjl (patched shipped: qjl.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/qjl.metal kernel_attn_score_qjl1_256 fixtures/qjl.json && \
|
||
echo "[verify-fork] metal/polar (patched shipped: polar.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/polar.metal kernel_mul_mv_q4_polar_f32 fixtures/polar.json && \
|
||
echo "[verify-fork] metal/polar_preht (patched shipped: polar.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/polar.metal kernel_mul_mv_q4_polar_preht_f32 fixtures/polar.json && \
|
||
echo "[verify-fork] metal/polar_preht + QJL residual (patched shipped: polar.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/polar.metal kernel_mul_mv_q4_polar_preht_f32 fixtures/polar_qjl.json && \
|
||
echo "[verify-fork] metal/fused_attn_qjl_tbq (patched shipped: fused_attn_qjl_tbq.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/fused_attn_qjl_tbq.metal kernel_fused_attn_qjl_tbq3_f32 fixtures/fused_attn_qjl_tbq.json && \
|
||
echo "[verify-fork] metal/fused_attn_qjl_tbq causal prefix (patched shipped: fused_attn_qjl_tbq.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/fused_attn_qjl_tbq.metal kernel_fused_attn_qjl_tbq3_f32 fixtures/fused_attn_qjl_tbq_causal.json && \
|
||
echo "[verify-fork] metal/fused_attn_qjl_polar (patched shipped: fused_attn_qjl_polar.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/fused_attn_qjl_polar.metal kernel_fused_attn_qjl_polar_f32 fixtures/fused_attn_qjl_polar.json && \
|
||
echo "[verify-fork] metal/fused_attn_qjl_polar causal prefix (patched shipped: fused_attn_qjl_polar.metal)" && ./metal_verify $(ELIZA_MTP_FORK_METAL_DIR)/fused_attn_qjl_polar.metal kernel_fused_attn_qjl_polar_f32 fixtures/fused_attn_qjl_polar_causal.json && \
|
||
echo "[verify-fork] metal: 7/7 patched shipped shader files plus Polar preht/fused-attn variants verified against the same fixtures as the standalones"; \
|
||
else \
|
||
echo "[verify-fork] xcrun not found — skipping Metal fork verification (macOS only)"; \
|
||
fi
|
||
@if [ -d "$(ELIZA_MTP_FORK_VULKAN_DIR)" ] && ls $(ELIZA_MTP_FORK_VULKAN_DIR)/turbo3.comp >/dev/null 2>&1; then \
|
||
if [ -z "$(GLSLC)" ]; then \
|
||
echo "[verify-fork] vulkan: glslc not found — cannot compile fork .comp -> .spv; skipping"; \
|
||
elif $(MAKE) vulkan_verify >/dev/null 2>&1 && [ -x ./vulkan_verify ]; then \
|
||
mkdir -p bench_results/fork-spv; \
|
||
echo "[verify-fork] vulkan: compiling fork .comp -> .spv with $(GLSLC)"; \
|
||
for f in turbo3 turbo4 turbo3_tcq qjl polar; do \
|
||
$(GLSLC) --target-env=vulkan1.1 --target-spv=spv1.3 -fshader-stage=compute \
|
||
$(ELIZA_MTP_FORK_VULKAN_DIR)/$$f.comp -o bench_results/fork-spv/$$f.spv || exit 1; \
|
||
echo "[verify-fork] vulkan/$$f (fork: vulkan-shaders/$$f.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/$$f.spv fixtures/$$f.json || exit 1; \
|
||
done; \
|
||
echo "[verify-fork] vulkan/polar + QJL residual (fork: vulkan-shaders/polar.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/polar.spv fixtures/polar_qjl.json || exit 1; \
|
||
$(GLSLC) --target-env=vulkan1.1 --target-spv=spv1.3 -fshader-stage=compute \
|
||
$(ELIZA_MTP_FORK_VULKAN_DIR)/polar_preht.comp -o bench_results/fork-spv/polar_preht.spv || exit 1; \
|
||
echo "[verify-fork] vulkan/polar_preht (fork: vulkan-shaders/polar_preht.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/polar_preht.spv fixtures/polar.json || exit 1; \
|
||
echo "[verify-fork] vulkan/polar_preht + QJL residual (fork: vulkan-shaders/polar_preht.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/polar_preht.spv fixtures/polar_qjl.json || exit 1; \
|
||
for f in fused_attn_qjl_tbq fused_attn_qjl_polar; do \
|
||
$(GLSLC) --target-env=vulkan1.1 --target-spv=spv1.3 -fshader-stage=compute \
|
||
$(ELIZA_MTP_FORK_VULKAN_DIR)/$$f.comp -o bench_results/fork-spv/$$f.spv || exit 1; \
|
||
echo "[verify-fork] vulkan/$$f (fork: vulkan-shaders/$$f.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/$$f.spv fixtures/$$f.json || exit 1; \
|
||
done; \
|
||
echo "[verify-fork] vulkan/fused_attn_qjl_tbq causal prefix (fork: vulkan-shaders/fused_attn_qjl_tbq.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/fused_attn_qjl_tbq.spv fixtures/fused_attn_qjl_tbq_causal.json || exit 1; \
|
||
echo "[verify-fork] vulkan/fused_attn_qjl_polar causal prefix (fork: vulkan-shaders/fused_attn_qjl_polar.comp)"; \
|
||
./vulkan_verify bench_results/fork-spv/fused_attn_qjl_polar.spv fixtures/fused_attn_qjl_polar_causal.json || exit 1; \
|
||
echo "[verify-fork] vulkan: 8/8 fork score/fused shaders + polar residual/preht variants verified against standalone fixtures"; \
|
||
else \
|
||
echo "[verify-fork] vulkan_verify not buildable on this host — skipping Vulkan fork verification"; \
|
||
fi; \
|
||
else \
|
||
echo "[verify-fork] $(ELIZA_MTP_FORK_VULKAN_DIR) not staged with eliza shaders (run a *-vulkan build target first); skipping Vulkan fork verification"; \
|
||
fi
|
||
@echo "[verify-fork] done. Compare the per-kernel pass/fail above against the standalone reference runs to detect drift."
|
||
endif
|
||
|
||
# GGML_OP_ATTN_SCORE_QJL Metal dispatch smoke test. Links against the BUILT
|
||
# libggml-metal.dylib / libggml-base.dylib in the fork's
|
||
# build/darwin-arm64-metal/bin tree, so it exercises the actual fork dispatch
|
||
# pipeline (not the JIT). Prerequisite:
|
||
# `node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target
|
||
# darwin-arm64-metal` has succeeded and produced libggml-*.dylib + the
|
||
# eliza kernels in default.metallib.
|
||
ELIZA_MTP_METAL_DEFAULT_BUILD_DIR := $(ELIZA_MTP_LLAMA_DIR)/build/darwin-arm64-metal
|
||
ELIZA_MTP_METAL_FUSED_BUILD_DIR := $(ELIZA_MTP_LLAMA_DIR)/build/darwin-arm64-metal-fused
|
||
ELIZA_MTP_METAL_BUILD_DIR ?= $(if $(wildcard $(ELIZA_MTP_METAL_DEFAULT_BUILD_DIR)/bin/libggml-metal.dylib),$(ELIZA_MTP_METAL_DEFAULT_BUILD_DIR),$(if $(wildcard $(ELIZA_MTP_METAL_FUSED_BUILD_DIR)/bin/libggml-metal.dylib),$(ELIZA_MTP_METAL_FUSED_BUILD_DIR),$(ELIZA_MTP_METAL_DEFAULT_BUILD_DIR)))
|
||
ELIZA_MTP_METAL_BIN_DIR ?= $(ELIZA_MTP_METAL_BUILD_DIR)/bin
|
||
ELIZA_STATE_DIR ?= $(HOME)/.eliza
|
||
ELIZA_MTP_TARGET ?= linux-x64-vulkan
|
||
ifneq ($(strip $(ELIZA_MTP_TARGET_OUT_DIR)),)
|
||
ELIZA_MTP_VULKAN_OUT_DIR ?= $(ELIZA_MTP_TARGET_OUT_DIR)
|
||
else
|
||
ELIZA_MTP_VULKAN_OUT_DIR ?= $(ELIZA_STATE_DIR)/local-inference/bin/mtp/$(ELIZA_MTP_TARGET)
|
||
endif
|
||
ELIZA_MTP_VULKAN_BUILD_DIR ?= $(ELIZA_MTP_LLAMA_DIR)/build/$(ELIZA_MTP_TARGET)
|
||
ELIZA_MTP_VULKAN_BUILD_BIN_DIR := $(ELIZA_MTP_VULKAN_BUILD_DIR)/bin
|
||
ifndef ELIZA_MTP_VULKAN_BIN_DIR
|
||
ifneq ($(strip $(wildcard $(ELIZA_MTP_VULKAN_OUT_DIR)/libggml-vulkan.so)$(wildcard $(ELIZA_MTP_VULKAN_OUT_DIR)/libggml-vulkan.dylib)),)
|
||
ELIZA_MTP_VULKAN_BIN_DIR := $(ELIZA_MTP_VULKAN_OUT_DIR)
|
||
else
|
||
ELIZA_MTP_VULKAN_BIN_DIR := $(ELIZA_MTP_VULKAN_BUILD_BIN_DIR)
|
||
endif
|
||
endif
|
||
|
||
DISPATCH_SMOKE_GGML_HEADERS := $(wildcard \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml.h \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml-alloc.h \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml-backend.h \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml-metal.h \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml-vulkan.h \
|
||
$(ELIZA_MTP_LLAMA_DIR)/ggml/src/ggml-common.h)
|
||
DISPATCH_SMOKE_REF_HEADERS := ../reference/turbo_kernels.h qjl_polar_ref.h
|
||
|
||
dispatch-smoke: dispatch_smoke
|
||
@echo "[dispatch-smoke] running against $(ELIZA_MTP_METAL_BIN_DIR)"
|
||
GGML_METAL_PATH_RESOURCES=$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_METAL_BIN_DIR):$$DYLD_LIBRARY_PATH \
|
||
./dispatch_smoke
|
||
|
||
dispatch-smoke-implemented: dispatch_smoke
|
||
@echo "[dispatch-smoke] running implemented routes against $(ELIZA_MTP_METAL_BIN_DIR)"
|
||
GGML_METAL_PATH_RESOURCES=$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_METAL_BIN_DIR):$$DYLD_LIBRARY_PATH \
|
||
./dispatch_smoke
|
||
|
||
metal-upscale-probe: metal_upscale_probe
|
||
@echo "[metal-upscale-probe] probing CLIP UPSCALE fallback against $(ELIZA_MTP_METAL_BIN_DIR)"
|
||
GGML_METAL_PATH_RESOURCES=$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_METAL_BIN_DIR):$$DYLD_LIBRARY_PATH \
|
||
./metal_upscale_probe \
|
||
--out ../reports/porting/2026-05-12/vision-upscale-metal-probe.json
|
||
|
||
metal-upscale-probe-require: metal_upscale_probe
|
||
@echo "[metal-upscale-probe] requiring Metal support for CLIP UPSCALE"
|
||
GGML_METAL_PATH_RESOURCES=$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_METAL_BIN_DIR):$$DYLD_LIBRARY_PATH \
|
||
./metal_upscale_probe --require-metal-support
|
||
|
||
dispatch_smoke: dispatch_smoke.mm $(REF_OBJ) $(QJL_POLAR_OBJ) $(DISPATCH_SMOKE_GGML_HEADERS) $(DISPATCH_SMOKE_REF_HEADERS)
|
||
@which xcrun >/dev/null 2>&1 || { echo "xcrun not found — Metal dispatch-smoke needs macOS Xcode tools."; exit 1; }
|
||
@if [ ! -f "$(ELIZA_MTP_METAL_BIN_DIR)/libggml-metal.dylib" ]; then \
|
||
echo "[dispatch-smoke] missing $(ELIZA_MTP_METAL_BIN_DIR)/libggml-metal.dylib"; \
|
||
echo "[dispatch-smoke] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target darwin-arm64-metal' first."; \
|
||
exit 1; \
|
||
fi
|
||
clang++ $(CXXFLAGS) -ObjC++ -fobjc-arc \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/include \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/src \
|
||
-I../reference -I. \
|
||
dispatch_smoke.mm \
|
||
$(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
-L$(ELIZA_MTP_METAL_BIN_DIR) -lggml-base -lggml-metal -lggml-cpu -lggml \
|
||
-framework Foundation -framework Metal -framework MetalKit -framework CoreGraphics \
|
||
-Wl,-rpath,$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
-lm -o dispatch_smoke
|
||
|
||
metal_upscale_probe: metal_upscale_probe.cpp
|
||
@which xcrun >/dev/null 2>&1 || { echo "xcrun not found — Metal upscale probe needs macOS Xcode tools."; exit 1; }
|
||
@if [ ! -f "$(ELIZA_MTP_METAL_BIN_DIR)/libggml-metal.dylib" ]; then \
|
||
echo "[metal-upscale-probe] missing $(ELIZA_MTP_METAL_BIN_DIR)/libggml-metal.dylib"; \
|
||
echo "[metal-upscale-probe] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target darwin-arm64-metal' first."; \
|
||
exit 1; \
|
||
fi
|
||
clang++ $(CXXFLAGS) \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/include \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/src \
|
||
metal_upscale_probe.cpp \
|
||
-L$(ELIZA_MTP_METAL_BIN_DIR) -lggml-base -lggml-metal -lggml-cpu -lggml \
|
||
-framework Foundation -framework Metal -framework MetalKit -framework CoreGraphics \
|
||
-Wl,-rpath,$(ELIZA_MTP_METAL_BIN_DIR) \
|
||
-lm -o metal_upscale_probe
|
||
|
||
vulkan-dispatch-smoke: vulkan_dispatch_smoke
|
||
@if [ "$$(uname -s)" != "Linux" ] && [ "$${ELIZA_ALLOW_MOLTENVK_VULKAN_DISPATCH_SMOKE:-0}" != "1" ]; then \
|
||
echo "[vulkan-dispatch-smoke] native Linux required for recordable Vulkan graph dispatch evidence."; \
|
||
echo "[vulkan-dispatch-smoke] macOS/MoltenVK can validate standalone SPIR-V only; it cannot flip runtime-ready Vulkan capability bits."; \
|
||
echo "[vulkan-dispatch-smoke] run './linux_vulkan_smoke.sh' on physical Linux Vulkan hardware or use './android_vulkan_smoke.sh' with Android graph evidence."; \
|
||
exit 2; \
|
||
fi
|
||
@echo "[vulkan-dispatch-smoke] running against $(ELIZA_MTP_VULKAN_BIN_DIR)"
|
||
@echo "[vulkan-dispatch-smoke] managed output candidate: $(ELIZA_MTP_VULKAN_OUT_DIR)"
|
||
@echo "[vulkan-dispatch-smoke] build-tree candidate: $(ELIZA_MTP_VULKAN_BUILD_BIN_DIR)"
|
||
@echo "[vulkan-dispatch-smoke] requires ggml-vulkan to advertise and execute QJL/TBQ/Polar attention-score graph routes; symbol-only staging cannot pass"
|
||
@if [ ! -f "$(ELIZA_MTP_VULKAN_BIN_DIR)/libggml-vulkan.so" ] && [ ! -f "$(ELIZA_MTP_VULKAN_BIN_DIR)/libggml-vulkan.dylib" ]; then \
|
||
echo "[vulkan-dispatch-smoke] missing libggml-vulkan under $(ELIZA_MTP_VULKAN_BIN_DIR)"; \
|
||
echo "[vulkan-dispatch-smoke] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target $(ELIZA_MTP_TARGET)' on a Linux Vulkan host first."; \
|
||
exit 1; \
|
||
fi
|
||
LD_LIBRARY_PATH=$(ELIZA_MTP_VULKAN_BIN_DIR):$$LD_LIBRARY_PATH \
|
||
DYLD_LIBRARY_PATH=$(ELIZA_MTP_VULKAN_BIN_DIR):$$DYLD_LIBRARY_PATH \
|
||
./vulkan_dispatch_smoke
|
||
|
||
vulkan_dispatch_smoke: force vulkan_dispatch_smoke.cpp $(REF_OBJ) $(QJL_POLAR_OBJ) $(DISPATCH_SMOKE_GGML_HEADERS) $(DISPATCH_SMOKE_REF_HEADERS)
|
||
@if [ "$$(uname -s)" != "Linux" ] && [ "$${ELIZA_ALLOW_MOLTENVK_VULKAN_DISPATCH_SMOKE:-0}" != "1" ]; then \
|
||
echo "[vulkan-dispatch-smoke] native Linux required for recordable Vulkan graph dispatch evidence."; \
|
||
echo "[vulkan-dispatch-smoke] macOS/MoltenVK can validate standalone SPIR-V only; it cannot flip runtime-ready Vulkan capability bits."; \
|
||
echo "[vulkan-dispatch-smoke] run './linux_vulkan_smoke.sh' on physical Linux Vulkan hardware or use './android_vulkan_smoke.sh' with Android graph evidence."; \
|
||
exit 2; \
|
||
fi
|
||
@if [ ! -f "$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml.h" ] || [ ! -f "$(ELIZA_MTP_LLAMA_DIR)/ggml/include/ggml-backend.h" ]; then \
|
||
echo "[vulkan-dispatch-smoke] missing ggml headers under $(ELIZA_MTP_LLAMA_DIR)"; \
|
||
echo "[vulkan-dispatch-smoke] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target $(ELIZA_MTP_TARGET)' first, or override ELIZA_MTP_LLAMA_DIR."; \
|
||
exit 1; \
|
||
fi
|
||
@if [ ! -f "$(ELIZA_MTP_VULKAN_BIN_DIR)/libggml-vulkan.so" ] && [ ! -f "$(ELIZA_MTP_VULKAN_BIN_DIR)/libggml-vulkan.dylib" ]; then \
|
||
echo "[vulkan-dispatch-smoke] missing libggml-vulkan under $(ELIZA_MTP_VULKAN_BIN_DIR)"; \
|
||
echo "[vulkan-dispatch-smoke] managed output candidate: $(ELIZA_MTP_VULKAN_OUT_DIR)"; \
|
||
echo "[vulkan-dispatch-smoke] build-tree candidate: $(ELIZA_MTP_VULKAN_BUILD_BIN_DIR)"; \
|
||
echo "[vulkan-dispatch-smoke] run 'node packages/app-core/scripts/build-llama-cpp-mtp.mjs --target $(ELIZA_MTP_TARGET)' on a Linux Vulkan host first."; \
|
||
echo "[vulkan-dispatch-smoke] if using explicit prebuilt artifacts, set ELIZA_MTP_VULKAN_BIN_DIR=/path/to/the/bin-dir-containing-libggml-vulkan.so."; \
|
||
exit 1; \
|
||
fi
|
||
$(CXX) $(CXXFLAGS) \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/include \
|
||
-I$(ELIZA_MTP_LLAMA_DIR)/ggml/src \
|
||
-I../reference -I. \
|
||
vulkan_dispatch_smoke.cpp \
|
||
$(REF_OBJ) $(QJL_POLAR_OBJ) \
|
||
-L$(ELIZA_MTP_VULKAN_BIN_DIR) -lggml-base -lggml-vulkan -lggml-cpu -lggml \
|
||
-Wl,-rpath,$(ELIZA_MTP_VULKAN_BIN_DIR) \
|
||
-lm -o vulkan_dispatch_smoke
|
||
|
||
force:
|
||
|
||
clean:
|
||
rm -f $(REF_OBJ) $(QJL_POLAR_OBJ) gen_fixture vulkan_verify metal_verify metal_bench metal_upscale_probe cpu_bench cpu_simd_bench cuda_verify dispatch_smoke vulkan_dispatch_smoke
|
||
rm -f ../vulkan/*.spv
|