Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
28 commits
Select commit Hold shift + click to select a range
33b4653
Add cross-platform packed inference paths
iamwavecut Jul 11, 2026
5ea2e26
Relax realistic AVX2 BF16 test tolerance
iamwavecut Jul 11, 2026
fe6bbf4
ci: compile Vulkan sources on Windows
iamwavecut Jul 12, 2026
3c6b3d7
ci: discover the installed MSVC toolchain
iamwavecut Jul 12, 2026
f49f17c
ci: use Ninja with the discovered MSVC
iamwavecut Jul 12, 2026
667d216
ci: select the installed Visual Studio generator
iamwavecut Jul 12, 2026
e703453
ci: normalize ExecuTorch sources for Windows codegen
iamwavecut Jul 12, 2026
f41ae5e
Fix Vulkan PTE build with CMake 4
iamwavecut Jul 12, 2026
5ac63f8
Document Windows Vulkan build evidence
iamwavecut Jul 12, 2026
3afdc7e
ci: package Windows Vulkan runners
iamwavecut Jul 12, 2026
9a1ffea
ci: record Windows bundle provenance
iamwavecut Jul 12, 2026
85d2d28
ci: require the official MSVC runtime
iamwavecut Jul 12, 2026
6264726
ci: bundle transitive Windows notices
iamwavecut Jul 12, 2026
faac9d8
ci: tighten Windows bundle license scope
iamwavecut Jul 12, 2026
834474c
ci: cap the complete Windows bundle
iamwavecut Jul 12, 2026
30e95a4
Fix CUDA W4A4 and AdaLN paths; speed up Triton activation and matmul
iamwavecut Jul 12, 2026
2cad1c0
Add SIMD decoders for W2/W3/W6 packed CPU matmul
iamwavecut Jul 12, 2026
4932834
Reuse a persistent CPU worker pool and cache activation-constant vali…
iamwavecut Jul 12, 2026
a7667ec
Document CPU low-bit SIMD coverage and the worker pool in the kernel …
iamwavecut Jul 12, 2026
458f5fc
Tune CPU matmul by CPU family and buffer per-column decode; vectorize…
iamwavecut Jul 12, 2026
452607f
Add a cp.async-pipelined mma64 kernel, enabled by default for W2 only
iamwavecut Jul 12, 2026
0317f3b
Size MPS threadgroup memory per RPBH block and simplify reductions
iamwavecut Jul 12, 2026
3713a02
Fuse the W4A4 GEMM in Triton, cache forward-plan state, and support t…
iamwavecut Jul 12, 2026
a4a526a
Split the CLI/hub test monoliths, reconcile audit docs, and ship py.t…
iamwavecut Jul 12, 2026
b8cb1a5
Refresh the CPU low-bit benchmark table and bump to 0.5.0
iamwavecut Jul 12, 2026
1c2ef11
Gate the mma64 cp.async pipeline by grid occupancy and harden multi-a…
iamwavecut Jul 12, 2026
3d4208a
Tune fused W4A4 tiles by batch size and raise the int_mm crossover
iamwavecut Jul 12, 2026
b8ab9e8
Ship fused W2A4 and direct INT8-activation GEMM paths
iamwavecut Jul 12, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
269 changes: 269 additions & 0 deletions .github/workflows/windows-vulkan.yml
Original file line number Diff line number Diff line change
@@ -0,0 +1,269 @@
name: Windows Vulkan build

on:
push:
branches: ["main"]
paths:
- ".github/workflows/windows-vulkan.yml"
- "native-kernels/orbitquant-vulkan/**"
pull_request:
paths:
- ".github/workflows/windows-vulkan.yml"
- "native-kernels/orbitquant-vulkan/**"
workflow_dispatch:

permissions:
contents: read

concurrency:
group: windows-vulkan-${{ github.event.pull_request.number || github.ref }}
cancel-in-progress: true

env:
EXECUTORCH_COMMIT: a6d812a082df57898b8608f56c867140cc9da32c
ORBITQUANT_SOURCE_SHA: ${{ github.event.pull_request.head.sha || github.sha }}

jobs:
build:
name: MSVC Vulkan source build
runs-on: windows-2025
timeout-minutes: 90
steps:
- name: Check out OrbitQuant
uses: actions/checkout@v7.0.0

- name: Set up Python
uses: actions/setup-python@v6.3.0
with:
python-version: "3.12"

- name: Build the ExecuTorch and OrbitQuant Vulkan targets
shell: pwsh
run: |
$ErrorActionPreference = "Stop"
$PSNativeCommandUseErrorActionPreference = $true

$executorch = Join-Path $env:RUNNER_TEMP "executorch"
$executorchBuild = Join-Path $env:RUNNER_TEMP "executorch-build"
$executorchInstall = Join-Path $env:RUNNER_TEMP "executorch-install"
$orbitquantBuild = Join-Path $env:RUNNER_TEMP "orbitquant-vulkan-build"
$python = Join-Path $env:pythonLocation "python.exe"

git config --global core.autocrlf false
git config --global core.symlinks true
git clone --filter=blob:none --no-checkout `
https://github.com/pytorch/executorch.git $executorch
git -C $executorch fetch --depth=1 origin $env:EXECUTORCH_COMMIT
git -C $executorch config core.autocrlf false
git -C $executorch config core.symlinks true
git -C $executorch checkout --detach $env:EXECUTORCH_COMMIT
if ((git -C $executorch rev-parse HEAD) -ne $env:EXECUTORCH_COMMIT) {
throw "ExecuTorch checkout does not match the pinned commit"
}
git -C $executorch submodule update --init `
backends/vulkan/third-party/Vulkan-Headers `
backends/vulkan/third-party/VulkanMemoryAllocator `
backends/vulkan/third-party/volk
git -C $executorch submodule update --init

$env:PATH = `
"$env:CONDA;$env:CONDA\Scripts;$env:CONDA\Library\bin;$env:PATH"
Push-Location $executorch
& .ci\scripts\setup-vulkan-windows-deps.ps1
Pop-Location
$glslc = (Get-Command glslc.exe -ErrorAction Stop).Source
& $python -m pip install PyYAML==6.0.1
& $python -m pip install torch==2.12.1
$env:PYTHONPATH = `
"$executorch\src;$executorch;$(Split-Path $executorch -Parent)"
& $python -c `
"import executorch.codegen, torchgen; print('codegen imports ok')"

$vswhere = `
"${env:ProgramFiles(x86)}\Microsoft Visual Studio\Installer\vswhere.exe"
$vsRoot = & $vswhere `
-latest `
-products * `
-requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 `
-property installationPath
if (-not $vsRoot) {
throw "MSVC x86_64 tools were not found"
}
$vsVersion = & $vswhere `
-latest `
-products * `
-requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 `
-property installationVersion
$vsMajor = [int](($vsVersion -split "\.")[0])
$cmakeGenerator = switch ($vsMajor) {
17 { "Visual Studio 17 2022" }
18 { "Visual Studio 18 2026" }
default { throw "Unsupported Visual Studio major version $vsMajor" }
}
$vsDevCmd = Join-Path $vsRoot "Common7\Tools\VsDevCmd.bat"
$envDump = cmd.exe /d /s /c `
"call `"$vsDevCmd`" -arch=x64 -host_arch=x64 >nul && set"
foreach ($line in $envDump) {
$parts = $line -split "=", 2
if ($parts.Length -eq 2) {
Set-Item -Path "Env:$($parts[0])" -Value $parts[1]
}
}
Get-Command cl.exe -ErrorAction Stop | Format-List Source
Write-Host "Using CMake generator: $cmakeGenerator"
$glslcVersion = (& $glslc --version | Select-Object -First 1)
Write-Host $glslcVersion
cmake --version

cmake -S $executorch -B $executorchBuild `
-G $cmakeGenerator -A x64 `
-DCMAKE_BUILD_TYPE=Release `
"-DCMAKE_INSTALL_PREFIX=$executorchInstall" `
"-DPYTHON_EXECUTABLE=$python" `
-DEXECUTORCH_BUILD_VULKAN=ON `
-DEXECUTORCH_BUILD_PORTABLE_OPS=ON `
-DEXECUTORCH_BUILD_EXECUTOR_RUNNER=OFF `
-DEXECUTORCH_BUILD_TESTS=OFF `
-DEXECUTORCH_BUILD_PYBIND=OFF `
-DEXECUTORCH_BUILD_XNNPACK=OFF `
"-DGLSLC_PATH=$glslc"
cmake --build $executorchBuild --config Release `
--target install --parallel 4

cmake -S native-kernels\orbitquant-vulkan -B $orbitquantBuild `
-G $cmakeGenerator -A x64 `
-DCMAKE_BUILD_TYPE=Release `
"-DCMAKE_PREFIX_PATH=$executorchInstall" `
"-DPYTHON_EXECUTABLE=$python" `
"-DORBITQUANT_EXECUTORCH_ROOT=$executorch" `
"-DGLSLC_PATH=$glslc" `
-DORBITQUANT_VULKAN_BUILD_TESTS=ON `
-DORBITQUANT_VULKAN_BUILD_PTE_RUNNER=ON
cmake --build $orbitquantBuild --config Release `
--target orbitquant_vulkan_test orbitquant_vulkan_pte_runner `
--parallel 4

Get-Item `
"$orbitquantBuild\Release\orbitquant_vulkan_test.exe", `
"$orbitquantBuild\Release\orbitquant_vulkan_pte_runner.exe"

$bundle = Join-Path `
$env:RUNNER_TEMP "orbitquant-vulkan-windows-x86_64"
New-Item -ItemType Directory -Force -Path $bundle | Out-Null
Copy-Item `
"$orbitquantBuild\Release\orbitquant_vulkan_test.exe", `
"$orbitquantBuild\Release\orbitquant_vulkan_pte_runner.exe" `
-Destination $bundle

$runtimeDirs = @(
"$orbitquantBuild\Release",
"$executorchInstall\bin"
)
foreach ($runtimeDir in $runtimeDirs) {
if (Test-Path $runtimeDir) {
Get-ChildItem $runtimeDir -Filter "*.dll" -File |
Copy-Item -Destination $bundle -Force
}
}

$licenseFiles = @(
@{
Source = "$env:GITHUB_WORKSPACE\LICENSE"
Destination = "LICENSE-OrbitQuant.txt"
},
@{
Source = "$executorch\LICENSE"
Destination = "LICENSE-ExecuTorch.txt"
},
@{
Source = "$executorch\third-party\flatbuffers\LICENSE"
Destination = "LICENSE-FlatBuffers.txt"
},
@{
Source = "$executorch\third-party\gflags\COPYING.txt"
Destination = "LICENSE-gflags.txt"
},
@{
Source = "$executorch\backends\vulkan\third-party\Vulkan-Headers\LICENSE.md"
Destination = "LICENSE-Vulkan-Headers.md"
},
@{
Source = "$executorch\backends\vulkan\third-party\VulkanMemoryAllocator\LICENSE.txt"
Destination = "LICENSE-VulkanMemoryAllocator.txt"
},
@{
Source = "$executorch\backends\vulkan\third-party\volk\LICENSE.md"
Destination = "LICENSE-volk.md"
}
)
foreach ($licenseFile in $licenseFiles) {
if (-not (Test-Path $licenseFile.Source)) {
throw "Required license is missing: $($licenseFile.Source)"
}
Copy-Item `
$licenseFile.Source `
(Join-Path $bundle $licenseFile.Destination)
}

@(
"Install the current Microsoft Visual C++ v14 x64 Redistributable:",
"https://aka.ms/vc14/vc_redist.x64.exe",
"",
"Prerequisite: install the hardware vendor Vulkan driver.",
"The bundle intentionally does not ship vulkan-1.dll.",
"",
"Run native correctness and realistic-shape benchmarks:",
".\orbitquant_vulkan_test.exe",
"",
"Run an exported PTE with one raw tensor input:",
".\orbitquant_vulkan_pte_runner.exe --model_path model.pte --inputs input0.bin --output_file actual --print_output=none",
"The output tensor is written to actual-0.bin."
) | Set-Content `
-Encoding ascii (Join-Path $bundle "README.txt")

$dependencyReport = Join-Path $bundle "DEPENDENCIES.txt"
foreach ($executable in Get-ChildItem $bundle -Filter "*.exe") {
"[$($executable.Name)]" | Add-Content $dependencyReport
dumpbin.exe /dependents $executable.FullName |
Add-Content $dependencyReport
}
@(
"orbitquant_source_commit=$env:ORBITQUANT_SOURCE_SHA",
"orbitquant_checkout_commit=$(git rev-parse HEAD)",
"executorch_commit=$env:EXECUTORCH_COMMIT",
"cmake_generator=$cmakeGenerator",
"visual_studio_version=$vsVersion",
"python_version=3.12",
"pytorch_version=2.12.1",
"glslc_version=$glslcVersion"
) | Set-Content `
-Encoding ascii (Join-Path $bundle "BUILD-METADATA.txt")
$checksumPath = Join-Path $bundle "SHA256SUMS"
Get-ChildItem $bundle -File -Recurse |
Where-Object { $_.FullName -ne $checksumPath } |
Sort-Object FullName |
ForEach-Object {
$hash = (Get-FileHash $_.FullName -Algorithm SHA256).Hash
$relativePath = [IO.Path]::GetRelativePath(
$bundle, $_.FullName
).Replace("\", "/")
"$hash $relativePath"
} | Set-Content -Encoding ascii $checksumPath

$bundleBytes = (
Get-ChildItem $bundle -File -Recurse | Measure-Object Length -Sum
).Sum
if ($bundleBytes -gt 64MB) {
throw "Windows Vulkan runner bundle exceeds 64 MiB: $bundleBytes"
}
Write-Host "Windows Vulkan runner bundle bytes: $bundleBytes"
Get-ChildItem $bundle | Format-Table Name, Length

- name: Upload Windows Vulkan runner bundle
uses: actions/upload-artifact@v7.0.1
with:
name: orbitquant-vulkan-windows-x86_64-${{ env.ORBITQUANT_SOURCE_SHA }}
path: ${{ runner.temp }}/orbitquant-vulkan-windows-x86_64
if-no-files-found: error
retention-days: 14
compression-level: 9
79 changes: 79 additions & 0 deletions CLAUDE.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,79 @@
# CLAUDE.md

Guidance for coding agents working in this repository (OrbitQuant: calibration-free
post-training quantizer for transformer linear projections).

## Environment

```bash
uv sync --extra hf --extra dev
```

- Tests: `uv run pytest -ra`
- Lint: `uv run ruff check .` (line-length 100)

## Formatting

- NEVER mass-reformat. `ruff format .` is forbidden: the tree has historical
formatting drift and a blanket reformat produces unreviewable diffs.
- Format only the lines you touch.

## Gates Before Claiming Done

Run all of these and confirm they pass:

```bash
uv run ruff check .
uv run pytest -ra
scripts/run_paper_methodology_checks.sh
scripts/run_hf_compat_checks.sh --mode all
```

Backend gates (run on matching hardware only):

```bash
scripts/run_cuda_kernel_checks.sh # CUDA host
PYTHON_BIN="$(uv python find)" scripts/run_mps_kernel_checks.sh # Apple Silicon
```

The MPS gate requires bash >= 4; on macOS use `/opt/homebrew/bin/bash`.

## Native Kernels

The native package lives in `native-kernels/orbitquant-packed-matmul` and uses the
Hugging Face kernel-builder, pinned to rev
`d43de01d0b43285d8e5061ca4380c2bd1c40ae3b` (see `.github/workflows/ci.yml`).

Local build:

```bash
cd native-kernels/orbitquant-packed-matmul
kernel-builder create-pyproject -f .
CMAKE_ARGS=-DGPU_LANG=CPU python setup.py build_kernel
```

or via Nix:

```bash
nix --option sandbox relaxed run .#build-and-copy -L
```

Build artifacts inside the kernel directory are GENERATED — never commit them:
`CMakeLists.txt`, `cmake/`, `compat.py`, `metadata-*.json`, `pyproject.toml`,
`setup.py`, `build/`, `torch-ext/registration.h`,
`torch-ext/orbitquant_packed_matmul/_ops.py`, and `uv.lock` inside the kernel
directory. Delete them before running the repo test suite: two ABI3 lint tests
scan the kernel directory and fail on leftover generated files.

## Documentation Policy

- Docs describe the current artifact and user contract only. No development
chronicles, no raw logs.
- The source of truth for backend status is
`orbitquant.kernels.dispatch.backend_capabilities()`; the audited backend
contract is `docs/kernel-audit.md`.

## Hard Boundaries

- Hugging Face Kernel Hub publication is OUT OF SCOPE. Never post there.
- Do not commit, push, or publish anything without explicit operator instruction.
6 changes: 6 additions & 0 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,12 @@ The implementation is clean-room and Apache-2.0 licensed.
`from_pretrained()` integration.
- Packed-weight CUDA, Triton, and Metal inference paths that avoid a full
dequantized weight matrix.
- `torch.compile(fullgraph=True)` support: quantized forwards run behind a
registered custom op with a fake implementation, so compiled models match
eager output exactly. The forward is also CUDA-Graph-capturable;
`torch.compile(mode="reduce-overhead")` removes the per-launch overhead that
dominates small decode batches (measured 1.6-5.5x at 1-32 rows on
A40/RTX 4090).

Embeddings, timestep modules, task heads, and common final projections are
kept in source precision by default. Every automatic decision is available as
Expand Down
Loading